diff options
Diffstat (limited to 'clang/lib/Driver')
21 files changed, 544 insertions, 602 deletions
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index eeb4823..853f694 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -68,6 +68,7 @@ #include "clang/Driver/Types.h" #include "llvm/ADT/ArrayRef.h" #include "llvm/ADT/STLExtras.h" +#include "llvm/ADT/SmallSet.h" #include "llvm/ADT/StringExtras.h" #include "llvm/ADT/StringRef.h" #include "llvm/ADT/StringSet.h" @@ -83,6 +84,7 @@ #include "llvm/Support/ErrorHandling.h" #include "llvm/Support/ExitCodes.h" #include "llvm/Support/FileSystem.h" +#include "llvm/Support/FileUtilities.h" #include "llvm/Support/FormatVariadic.h" #include "llvm/Support/MD5.h" #include "llvm/Support/Path.h" @@ -109,65 +111,6 @@ using namespace clang::driver; using namespace clang; using namespace llvm::opt; -static std::optional<llvm::Triple> getOffloadTargetTriple(const Driver &D, - const ArgList &Args) { - auto OffloadTargets = Args.getAllArgValues(options::OPT_offload_EQ); - // Offload compilation flow does not support multiple targets for now. We - // need the HIPActionBuilder (and possibly the CudaActionBuilder{,Base}too) - // to support multiple tool chains first. - switch (OffloadTargets.size()) { - default: - D.Diag(diag::err_drv_only_one_offload_target_supported); - return std::nullopt; - case 0: - D.Diag(diag::err_drv_invalid_or_unsupported_offload_target) << ""; - return std::nullopt; - case 1: - break; - } - return llvm::Triple(OffloadTargets[0]); -} - -static std::optional<llvm::Triple> -getNVIDIAOffloadTargetTriple(const Driver &D, const ArgList &Args, - const llvm::Triple &HostTriple) { - if (!Args.hasArg(options::OPT_offload_EQ)) { - return llvm::Triple(HostTriple.isArch64Bit() ? "nvptx64-nvidia-cuda" - : "nvptx-nvidia-cuda"); - } - auto TT = getOffloadTargetTriple(D, Args); - if (TT && (TT->getArch() == llvm::Triple::spirv32 || - TT->getArch() == llvm::Triple::spirv64)) { - if (Args.hasArg(options::OPT_emit_llvm)) - return TT; - D.Diag(diag::err_drv_cuda_offload_only_emit_bc); - return std::nullopt; - } - D.Diag(diag::err_drv_invalid_or_unsupported_offload_target) << TT->str(); - return std::nullopt; -} - -static std::optional<llvm::Triple> -getHIPOffloadTargetTriple(const Driver &D, const ArgList &Args) { - if (!Args.hasArg(options::OPT_offload_EQ)) { - auto OffloadArchs = Args.getAllArgValues(options::OPT_offload_arch_EQ); - if (llvm::is_contained(OffloadArchs, "amdgcnspirv") && - OffloadArchs.size() == 1) - return llvm::Triple("spirv64-amd-amdhsa"); - return llvm::Triple("amdgcn-amd-amdhsa"); // Default HIP triple. - } - auto TT = getOffloadTargetTriple(D, Args); - if (!TT) - return std::nullopt; - if (TT->isAMDGCN() && TT->getVendor() == llvm::Triple::AMD && - TT->getOS() == llvm::Triple::AMDHSA) - return TT; - if (TT->getArch() == llvm::Triple::spirv64) - return TT; - D.Diag(diag::err_drv_invalid_or_unsupported_offload_target) << TT->str(); - return std::nullopt; -} - template <typename F> static bool usesInput(const ArgList &Args, F &&Fn) { return llvm::any_of(Args, [&](Arg *A) { return (A->getOption().matches(options::OPT_x) && @@ -266,8 +209,8 @@ Driver::Driver(StringRef ClangExecutable, StringRef TargetTriple, CCLogDiagnostics(false), CCGenDiagnostics(false), CCPrintProcessStats(false), CCPrintInternalStats(false), TargetTriple(TargetTriple), Saver(Alloc), PrependArg(nullptr), - CheckInputsExist(true), ProbePrecompiled(true), - SuppressMissingInputWarning(false) { + PreferredLinker(CLANG_DEFAULT_LINKER), CheckInputsExist(true), + ProbePrecompiled(true), SuppressMissingInputWarning(false) { // Provide a sane fallback if no VFS is specified. if (!this->VFS) this->VFS = llvm::vfs::getRealFileSystem(); @@ -458,6 +401,44 @@ phases::ID Driver::getFinalPhase(const DerivedArgList &DAL, return FinalPhase; } +llvm::Expected<std::unique_ptr<llvm::MemoryBuffer>> +Driver::executeProgram(llvm::ArrayRef<llvm::StringRef> Args) const { + llvm::SmallString<64> OutputFile; + llvm::sys::fs::createTemporaryFile("driver-program", "txt", OutputFile, + llvm::sys::fs::OF_Text); + llvm::FileRemover OutputRemover(OutputFile.c_str()); + std::optional<llvm::StringRef> Redirects[] = { + {""}, + OutputFile.str(), + {""}, + }; + + std::string ErrorMessage; + int SecondsToWait = 60; + if (std::optional<std::string> Str = + llvm::sys::Process::GetEnv("CLANG_TOOLCHAIN_PROGRAM_TIMEOUT")) { + if (!llvm::to_integer(*Str, SecondsToWait)) + return llvm::createStringError(std::error_code(), + "CLANG_TOOLCHAIN_PROGRAM_TIMEOUT expected " + "an integer, got '" + + *Str + "'"); + SecondsToWait = std::max(SecondsToWait, 0); // infinite + } + StringRef Executable = Args[0]; + if (llvm::sys::ExecuteAndWait(Executable, Args, {}, Redirects, SecondsToWait, + /*MemoryLimit=*/0, &ErrorMessage)) + return llvm::createStringError(std::error_code(), + Executable + ": " + ErrorMessage); + + llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> OutputBuf = + llvm::MemoryBuffer::getFile(OutputFile.c_str()); + if (!OutputBuf) + return llvm::createStringError(OutputBuf.getError(), + "Failed to read stdout of " + Executable + + ": " + OutputBuf.getError().message()); + return std::move(*OutputBuf); +} + static Arg *MakeInputArg(DerivedArgList &Args, const OptTable &Opts, StringRef Value, bool Claim = true) { Arg *A = new Arg(Opts.getOption(options::OPT_INPUT), Value, @@ -921,250 +902,266 @@ Driver::OpenMPRuntimeKind Driver::getOpenMPRuntime(const ArgList &Args) const { return RT; } -static llvm::Triple getSYCLDeviceTriple(StringRef TargetArch) { - SmallVector<StringRef, 5> SYCLAlias = {"spir", "spir64", "spirv", "spirv32", - "spirv64"}; - if (llvm::is_contained(SYCLAlias, TargetArch)) { - llvm::Triple TargetTriple; - TargetTriple.setArchName(TargetArch); - TargetTriple.setVendor(llvm::Triple::UnknownVendor); - TargetTriple.setOS(llvm::Triple::UnknownOS); - return TargetTriple; - } - return llvm::Triple(TargetArch); +// Handles `native` offload architectures by using the 'offload-arch' utility. +static llvm::SmallVector<std::string> +getSystemOffloadArchs(Compilation &C, Action::OffloadKind Kind) { + StringRef Program = C.getArgs().getLastArgValue( + options::OPT_offload_arch_tool_EQ, "offload-arch"); + + SmallVector<std::string> GPUArchs; + if (llvm::ErrorOr<std::string> Executable = + llvm::sys::findProgramByName(Program)) { + llvm::SmallVector<StringRef> Args{*Executable}; + if (Kind == Action::OFK_HIP) + Args.push_back("--only=amdgpu"); + else if (Kind == Action::OFK_Cuda) + Args.push_back("--only=nvptx"); + auto StdoutOrErr = C.getDriver().executeProgram(Args); + + if (!StdoutOrErr) { + C.getDriver().Diag(diag::err_drv_undetermined_gpu_arch) + << Action::GetOffloadKindName(Kind) << StdoutOrErr.takeError() + << "--offload-arch"; + return GPUArchs; + } + if ((*StdoutOrErr)->getBuffer().empty()) { + C.getDriver().Diag(diag::err_drv_undetermined_gpu_arch) + << Action::GetOffloadKindName(Kind) << "No GPU detected in the system" + << "--offload-arch"; + return GPUArchs; + } + + for (StringRef Arch : llvm::split((*StdoutOrErr)->getBuffer(), "\n")) + if (!Arch.empty()) + GPUArchs.push_back(Arch.str()); + } else { + C.getDriver().Diag(diag::err_drv_command_failure) << "offload-arch"; + } + return GPUArchs; } -static bool addSYCLDefaultTriple(Compilation &C, - SmallVectorImpl<llvm::Triple> &SYCLTriples) { - // Check current set of triples to see if the default has already been set. - for (const auto &SYCLTriple : SYCLTriples) { - if (SYCLTriple.getSubArch() == llvm::Triple::NoSubArch && - SYCLTriple.isSPIROrSPIRV()) - return false; +// Attempts to infer the correct offloading toolchain triple by looking at the +// requested offloading kind and architectures. +static llvm::DenseSet<llvm::StringRef> +inferOffloadToolchains(Compilation &C, Action::OffloadKind Kind) { + std::set<std::string> Archs; + for (Arg *A : C.getInputArgs()) { + for (StringRef Arch : A->getValues()) { + if (A->getOption().matches(options::OPT_offload_arch_EQ)) { + if (Arch == "native") { + for (StringRef Str : getSystemOffloadArchs(C, Kind)) + Archs.insert(Str.str()); + } else { + Archs.insert(Arch.str()); + } + } else if (A->getOption().matches(options::OPT_no_offload_arch_EQ)) { + if (Arch == "all") + Archs.clear(); + else + Archs.erase(Arch.str()); + } + } } - // Add the default triple as it was not found. - llvm::Triple DefaultTriple = getSYCLDeviceTriple( - C.getDefaultToolChain().getTriple().isArch32Bit() ? "spirv32" - : "spirv64"); - SYCLTriples.insert(SYCLTriples.begin(), DefaultTriple); - return true; + + llvm::DenseSet<llvm::StringRef> Triples; + for (llvm::StringRef Arch : Archs) { + OffloadArch ID = StringToOffloadArch(Arch); + if (ID == OffloadArch::UNKNOWN) + ID = StringToOffloadArch( + getProcessorFromTargetID(llvm::Triple("amdgcn-amd-amdhsa"), Arch)); + + if (Kind == Action::OFK_HIP && !IsAMDOffloadArch(ID)) { + C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch) + << "HIP" << Arch; + return llvm::DenseSet<llvm::StringRef>(); + } + if (Kind == Action::OFK_Cuda && !IsNVIDIAOffloadArch(ID)) { + C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch) + << "CUDA" << Arch; + return llvm::DenseSet<llvm::StringRef>(); + } + if (Kind == Action::OFK_OpenMP && + (ID == OffloadArch::UNKNOWN || ID == OffloadArch::UNUSED)) { + C.getDriver().Diag(clang::diag::err_drv_failed_to_deduce_target_from_arch) + << Arch; + return llvm::DenseSet<llvm::StringRef>(); + } + if (ID == OffloadArch::UNKNOWN || ID == OffloadArch::UNUSED) { + C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch) + << "offload" << Arch; + return llvm::DenseSet<llvm::StringRef>(); + } + + StringRef Triple; + if (ID == OffloadArch::AMDGCNSPIRV) + Triple = "spirv64-amd-amdhsa"; + else if (IsNVIDIAOffloadArch(ID)) + Triple = C.getDefaultToolChain().getTriple().isArch64Bit() + ? "nvptx64-nvidia-cuda" + : "nvptx-nvidia-cuda"; + else if (IsAMDOffloadArch(ID)) + Triple = "amdgcn-amd-amdhsa"; + else + continue; + + // Make a new argument that dispatches this argument to the appropriate + // toolchain. This is required when we infer it and create potentially + // incompatible toolchains from the global option. + Option Opt = C.getDriver().getOpts().getOption(options::OPT_Xarch__); + unsigned Index = C.getArgs().getBaseArgs().MakeIndex("-Xarch_"); + Arg *A = new Arg(Opt, C.getArgs().getArgString(Index), Index, + C.getArgs().MakeArgString(Triple.split("-").first), + C.getArgs().MakeArgString("--offload-arch=" + Arch)); + C.getArgs().append(A); + C.getArgs().AddSynthesizedArg(A); + Triples.insert(Triple); + } + + // Infer the default target triple if no specific architectures are given. + if (Archs.empty() && Kind == Action::OFK_HIP) + Triples.insert("amdgcn-amd-amdhsa"); + else if (Archs.empty() && Kind == Action::OFK_Cuda) + Triples.insert(C.getDefaultToolChain().getTriple().isArch64Bit() + ? "nvptx64-nvidia-cuda" + : "nvptx-nvidia-cuda"); + else if (Archs.empty() && Kind == Action::OFK_SYCL) + Triples.insert(C.getDefaultToolChain().getTriple().isArch64Bit() + ? "spirv64-unknown-unknown" + : "spirv32-unknown-unknown"); + + // We need to dispatch these to the appropriate toolchain now. + C.getArgs().eraseArg(options::OPT_offload_arch_EQ); + C.getArgs().eraseArg(options::OPT_no_offload_arch_EQ); + + return Triples; } void Driver::CreateOffloadingDeviceToolChains(Compilation &C, InputList &Inputs) { - - // - // CUDA/HIP - // - // We need to generate a CUDA/HIP toolchain if any of the inputs has a CUDA - // or HIP type. However, mixed CUDA/HIP compilation is not supported. + bool UseLLVMOffload = C.getInputArgs().hasArg( + options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false); bool IsCuda = - llvm::any_of(Inputs, [](std::pair<types::ID, const llvm::opt::Arg *> &I) { - return types::isCuda(I.first); - }); - bool IsHIP = llvm::any_of(Inputs, [](std::pair<types::ID, const llvm::opt::Arg *> &I) { - return types::isHIP(I.first); - }) || - C.getInputArgs().hasArg(options::OPT_hip_link) || - C.getInputArgs().hasArg(options::OPT_hipstdpar); - bool UseLLVMOffload = C.getInputArgs().hasArg( - options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false); - if (IsCuda && IsHIP) { - Diag(clang::diag::err_drv_mix_cuda_hip); + return types::isCuda(I.first); + }) && + !UseLLVMOffload; + bool IsHIP = + (llvm::any_of(Inputs, + [](std::pair<types::ID, const llvm::opt::Arg *> &I) { + return types::isHIP(I.first); + }) || + C.getInputArgs().hasArg(options::OPT_hip_link) || + C.getInputArgs().hasArg(options::OPT_hipstdpar)) && + !UseLLVMOffload; + bool IsSYCL = C.getInputArgs().hasFlag(options::OPT_fsycl, + options::OPT_fno_sycl, false); + bool IsOpenMPOffloading = + UseLLVMOffload || + (C.getInputArgs().hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ, + options::OPT_fno_openmp, false) && + (C.getInputArgs().hasArg(options::OPT_offload_targets_EQ) || + (C.getInputArgs().hasArg(options::OPT_offload_arch_EQ) && + !(IsCuda || IsHIP)))); + + llvm::SmallSet<Action::OffloadKind, 4> Kinds; + const std::pair<bool, Action::OffloadKind> ActiveKinds[] = { + {IsCuda, Action::OFK_Cuda}, + {IsHIP, Action::OFK_HIP}, + {IsOpenMPOffloading, Action::OFK_OpenMP}, + {IsSYCL, Action::OFK_SYCL}}; + for (const auto &[Active, Kind] : ActiveKinds) + if (Active) + Kinds.insert(Kind); + + // We currently don't support any kind of mixed offloading. + if (Kinds.size() > 1) { + Diag(clang::diag::err_drv_mix_offload) + << Action::GetOffloadKindName(*Kinds.begin()).upper() + << Action::GetOffloadKindName(*(++Kinds.begin())).upper(); return; } - if (IsCuda && !UseLLVMOffload) { - auto CudaTriple = getNVIDIAOffloadTargetTriple( - *this, C.getInputArgs(), C.getDefaultToolChain().getTriple()); - if (!CudaTriple) - return; - - auto &TC = - getOffloadToolChain(C.getInputArgs(), Action::OFK_Cuda, *CudaTriple, - C.getDefaultToolChain().getTriple()); - - // Emit a warning if the detected CUDA version is too new. - const CudaInstallationDetector &CudaInstallation = - static_cast<const toolchains::CudaToolChain &>(TC).CudaInstallation; - if (CudaInstallation.isValid()) - CudaInstallation.WarnIfUnsupportedVersion(); - C.addOffloadDeviceToolChain(&TC, Action::OFK_Cuda); - OffloadArchs[&TC] = getOffloadArchs(C, C.getArgs(), Action::OFK_Cuda, &TC, - /*SpecificToolchain=*/true); - } else if (IsHIP && !UseLLVMOffload) { - if (auto *OMPTargetArg = - C.getInputArgs().getLastArg(options::OPT_offload_targets_EQ)) { - Diag(clang::diag::err_drv_unsupported_opt_for_language_mode) - << OMPTargetArg->getSpelling() << "HIP"; - return; - } - - auto HIPTriple = getHIPOffloadTargetTriple(*this, C.getInputArgs()); - if (!HIPTriple) - return; - - auto &TC = - getOffloadToolChain(C.getInputArgs(), Action::OFK_HIP, *HIPTriple, - C.getDefaultToolChain().getTriple()); - C.addOffloadDeviceToolChain(&TC, Action::OFK_HIP); - - // TODO: Fix 'amdgcnspirv' handling with the new driver. - if (C.getInputArgs().hasFlag(options::OPT_offload_new_driver, - options::OPT_no_offload_new_driver, false)) - OffloadArchs[&TC] = getOffloadArchs(C, C.getArgs(), Action::OFK_HIP, &TC, - /*SpecificToolchain=*/true); - } + // Initialize the compilation identifier used for unique CUDA / HIP names. if (IsCuda || IsHIP) CUIDOpts = CUIDOptions(C.getArgs(), *this); - // - // OpenMP - // - // We need to generate an OpenMP toolchain if the user specified targets with - // the -fopenmp-targets option or used --offload-arch with OpenMP enabled. - bool IsOpenMPOffloading = - ((IsCuda || IsHIP) && UseLLVMOffload) || - (C.getInputArgs().hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ, - options::OPT_fno_openmp, false) && - (C.getInputArgs().hasArg(options::OPT_offload_targets_EQ) || - C.getInputArgs().hasArg(options::OPT_offload_arch_EQ))); - if (IsOpenMPOffloading) { - // We expect that -fopenmp-targets is always used in conjunction with the - // option -fopenmp specifying a valid runtime with offloading support, i.e. - // libomp or libiomp. - OpenMPRuntimeKind RuntimeKind = getOpenMPRuntime(C.getInputArgs()); - if (RuntimeKind != OMPRT_OMP && RuntimeKind != OMPRT_IOMP5) { - Diag(clang::diag::err_drv_expecting_fopenmp_with_fopenmp_targets); - return; - } - - // If the user specified -fopenmp-targets= we create a toolchain for each - // valid triple. Otherwise, if only --offload-arch= was specified we instead - // attempt to derive the appropriate toolchains from the arguments. - if (Arg *OpenMPTargets = - C.getInputArgs().getLastArg(options::OPT_offload_targets_EQ)) { - if (OpenMPTargets && !OpenMPTargets->getNumValues()) { - Diag(clang::diag::warn_drv_empty_joined_argument) - << OpenMPTargets->getAsString(C.getInputArgs()); + // Get the list of requested offloading toolchains. If they were not + // explicitly specified we will infer them based on the offloading language + // and requested architectures. + std::multiset<llvm::StringRef> Triples; + if (C.getInputArgs().hasArg(options::OPT_offload_targets_EQ)) { + std::vector<std::string> ArgValues = + C.getInputArgs().getAllArgValues(options::OPT_offload_targets_EQ); + for (llvm::StringRef Target : ArgValues) + Triples.insert(C.getInputArgs().MakeArgString(Target)); + + if (ArgValues.empty()) + Diag(clang::diag::warn_drv_empty_joined_argument) + << C.getInputArgs() + .getLastArg(options::OPT_offload_targets_EQ) + ->getAsString(C.getInputArgs()); + } else if (Kinds.size() > 0) { + for (Action::OffloadKind Kind : Kinds) { + llvm::DenseSet<llvm::StringRef> Derived = inferOffloadToolchains(C, Kind); + Triples.insert(Derived.begin(), Derived.end()); + } + } + + // Build an offloading toolchain for every requested target and kind. + llvm::StringMap<StringRef> FoundNormalizedTriples; + for (StringRef Target : Triples) { + // OpenMP offloading requires a compatible libomp. + if (Kinds.contains(Action::OFK_OpenMP)) { + OpenMPRuntimeKind RuntimeKind = getOpenMPRuntime(C.getInputArgs()); + if (RuntimeKind != OMPRT_OMP && RuntimeKind != OMPRT_IOMP5) { + Diag(clang::diag::err_drv_expecting_fopenmp_with_fopenmp_targets); return; } + } - // Make sure these show up in a deterministic order. - std::multiset<StringRef> OpenMPTriples; - for (StringRef T : OpenMPTargets->getValues()) - OpenMPTriples.insert(T); - - llvm::StringMap<StringRef> FoundNormalizedTriples; - for (StringRef T : OpenMPTriples) { - llvm::Triple TT(ToolChain::getOpenMPTriple(T)); - std::string NormalizedName = TT.normalize(); - - // Make sure we don't have a duplicate triple. - auto [TripleIt, Inserted] = - FoundNormalizedTriples.try_emplace(NormalizedName, T); - if (!Inserted) { - Diag(clang::diag::warn_drv_omp_offload_target_duplicate) - << T << TripleIt->second; - continue; - } - - // If the specified target is invalid, emit a diagnostic. - if (TT.getArch() == llvm::Triple::UnknownArch) { - Diag(clang::diag::err_drv_invalid_omp_target) << T; - continue; - } + // Certain options are not allowed when combined with SYCL compilation. + if (Kinds.contains(Action::OFK_SYCL)) { + for (auto ID : + {options::OPT_static_libstdcxx, options::OPT_ffreestanding}) + if (Arg *IncompatArg = C.getInputArgs().getLastArg(ID)) + Diag(clang::diag::err_drv_argument_not_allowed_with) + << IncompatArg->getSpelling() << "-fsycl"; + } - auto &TC = getOffloadToolChain(C.getInputArgs(), Action::OFK_OpenMP, TT, - C.getDefaultToolChain().getTriple()); - C.addOffloadDeviceToolChain(&TC, Action::OFK_OpenMP); - OffloadArchs[&TC] = - getOffloadArchs(C, C.getArgs(), Action::OFK_OpenMP, &TC, - /*SpecificToolchain=*/true); - } - } else if (C.getInputArgs().hasArg(options::OPT_offload_arch_EQ) && - ((!IsHIP && !IsCuda) || UseLLVMOffload)) { - llvm::Triple AMDTriple("amdgcn-amd-amdhsa"); - llvm::Triple NVPTXTriple("nvptx64-nvidia-cuda"); - - for (StringRef Arch : - C.getInputArgs().getAllArgValues(options::OPT_offload_arch_EQ)) { - bool IsNVPTX = IsNVIDIAOffloadArch( - StringToOffloadArch(getProcessorFromTargetID(NVPTXTriple, Arch))); - bool IsAMDGPU = IsAMDOffloadArch( - StringToOffloadArch(getProcessorFromTargetID(AMDTriple, Arch))); - if (!IsNVPTX && !IsAMDGPU && !Arch.empty() && - !Arch.equals_insensitive("native")) { - Diag(clang::diag::err_drv_failed_to_deduce_target_from_arch) << Arch; - return; - } + // Create a device toolchain for every specified kind and triple. + for (Action::OffloadKind Kind : Kinds) { + llvm::Triple TT = Kind == Action::OFK_OpenMP + ? ToolChain::getOpenMPTriple(Target) + : llvm::Triple(Target); + if (TT.getArch() == llvm::Triple::ArchType::UnknownArch) { + Diag(diag::err_drv_invalid_or_unsupported_offload_target) << TT.str(); + continue; } - // Attempt to deduce the offloading triple from the set of architectures. - // We can only correctly deduce NVPTX / AMDGPU triples currently. - for (const llvm::Triple &TT : {AMDTriple, NVPTXTriple}) { - auto &TC = getOffloadToolChain(C.getInputArgs(), Action::OFK_OpenMP, TT, - C.getDefaultToolChain().getTriple()); - - llvm::SmallVector<StringRef> Archs = - getOffloadArchs(C, C.getArgs(), Action::OFK_OpenMP, &TC, - /*SpecificToolchain=*/false); - if (!Archs.empty()) { - C.addOffloadDeviceToolChain(&TC, Action::OFK_OpenMP); - OffloadArchs[&TC] = Archs; - } + std::string NormalizedName = TT.normalize(); + auto [TripleIt, Inserted] = + FoundNormalizedTriples.try_emplace(NormalizedName, Target); + if (!Inserted) { + Diag(clang::diag::warn_drv_omp_offload_target_duplicate) + << Target << TripleIt->second; + continue; } - // If the set is empty then we failed to find a native architecture. - auto TCRange = C.getOffloadToolChains(Action::OFK_OpenMP); - if (TCRange.first == TCRange.second) - Diag(clang::diag::err_drv_failed_to_deduce_target_from_arch) - << "native"; - } - } else if (C.getInputArgs().hasArg(options::OPT_offload_targets_EQ)) { - Diag(clang::diag::err_drv_expecting_fopenmp_with_fopenmp_targets); - return; - } + auto &TC = getOffloadToolChain(C.getInputArgs(), Kind, TT, + C.getDefaultToolChain().getTriple()); - // We need to generate a SYCL toolchain if the user specified -fsycl. - bool IsSYCL = C.getInputArgs().hasFlag(options::OPT_fsycl, - options::OPT_fno_sycl, false); - - auto argSYCLIncompatible = [&](OptSpecifier OptId) { - if (!IsSYCL) - return; - if (Arg *IncompatArg = C.getInputArgs().getLastArg(OptId)) - Diag(clang::diag::err_drv_argument_not_allowed_with) - << IncompatArg->getSpelling() << "-fsycl"; - }; - // -static-libstdc++ is not compatible with -fsycl. - argSYCLIncompatible(options::OPT_static_libstdcxx); - // -ffreestanding cannot be used with -fsycl - argSYCLIncompatible(options::OPT_ffreestanding); - - llvm::SmallVector<llvm::Triple, 4> UniqueSYCLTriplesVec; - - if (IsSYCL) { - addSYCLDefaultTriple(C, UniqueSYCLTriplesVec); + // Emit a warning if the detected CUDA version is too new. + if (Kind == Action::OFK_Cuda) { + auto &CudaInstallation = + static_cast<const toolchains::CudaToolChain &>(TC).CudaInstallation; + if (CudaInstallation.isValid()) + CudaInstallation.WarnIfUnsupportedVersion(); + } - // We'll need to use the SYCL and host triples as the key into - // getOffloadingDeviceToolChain, because the device toolchains we're - // going to create will depend on both. - const ToolChain *HostTC = C.getSingleOffloadToolChain<Action::OFK_Host>(); - for (const auto &TT : UniqueSYCLTriplesVec) { - auto &TC = getOffloadToolChain(C.getInputArgs(), Action::OFK_SYCL, TT, - HostTC->getTriple()); - C.addOffloadDeviceToolChain(&TC, Action::OFK_SYCL); - OffloadArchs[&TC] = getOffloadArchs(C, C.getArgs(), Action::OFK_SYCL, &TC, - /*SpecificToolchain=*/true); + C.addOffloadDeviceToolChain(&TC, Kind); } } - - // - // TODO: Add support for other offloading programming models here. - // } bool Driver::loadZOSCustomizationFile(llvm::cl::ExpansionContext &ExpCtx) { @@ -3306,9 +3303,6 @@ class OffloadingActionBuilder final { // architecture. If we are in host-only mode we return 'success' so that // the host uses the CUDA offload kind. if (auto *IA = dyn_cast<InputAction>(HostAction)) { - assert(!GpuArchList.empty() && - "We should have at least one GPU architecture."); - // If the host input is not CUDA or HIP, we don't need to bother about // this input. if (!(IA->getType() == types::TY_CUDA || @@ -3408,10 +3402,6 @@ class OffloadingActionBuilder final { CudaDeviceActions.clear(); } - /// Get canonicalized offload arch option. \returns empty StringRef if the - /// option is invalid. - virtual StringRef getCanonicalOffloadArch(StringRef Arch) = 0; - virtual std::optional<std::pair<llvm::StringRef, llvm::StringRef>> getConflictOffloadArchCombination(const std::set<StringRef> &GpuArchs) = 0; @@ -3440,91 +3430,25 @@ class OffloadingActionBuilder final { return true; } - ToolChains.push_back( - AssociatedOffloadKind == Action::OFK_Cuda - ? C.getSingleOffloadToolChain<Action::OFK_Cuda>() - : C.getSingleOffloadToolChain<Action::OFK_HIP>()); - - CompileHostOnly = C.getDriver().offloadHostOnly(); - EmitLLVM = Args.getLastArg(options::OPT_emit_llvm); - EmitAsm = Args.getLastArg(options::OPT_S); - - // --offload and --offload-arch options are mutually exclusive. - if (Args.hasArgNoClaim(options::OPT_offload_EQ) && - Args.hasArgNoClaim(options::OPT_offload_arch_EQ, - options::OPT_no_offload_arch_EQ)) { - C.getDriver().Diag(diag::err_opt_not_valid_with_opt) << "--offload-arch" - << "--offload"; - } - - // Collect all offload arch parameters, removing duplicates. std::set<StringRef> GpuArchs; - bool Error = false; - const ToolChain &TC = *ToolChains.front(); - for (Arg *A : C.getArgsForToolChain(&TC, /*BoundArch=*/"", - AssociatedOffloadKind)) { - if (!(A->getOption().matches(options::OPT_offload_arch_EQ) || - A->getOption().matches(options::OPT_no_offload_arch_EQ))) - continue; - A->claim(); - - for (StringRef ArchStr : llvm::split(A->getValue(), ",")) { - if (A->getOption().matches(options::OPT_no_offload_arch_EQ) && - ArchStr == "all") { - GpuArchs.clear(); - } else if (ArchStr == "native") { - auto GPUsOrErr = ToolChains.front()->getSystemGPUArchs(Args); - if (!GPUsOrErr) { - TC.getDriver().Diag(diag::err_drv_undetermined_gpu_arch) - << llvm::Triple::getArchTypeName(TC.getArch()) - << llvm::toString(GPUsOrErr.takeError()) << "--offload-arch"; - continue; - } + for (Action::OffloadKind Kind : {Action::OFK_Cuda, Action::OFK_HIP}) { + for (auto &I : llvm::make_range(C.getOffloadToolChains(Kind))) { + ToolChains.push_back(I.second); - for (auto GPU : *GPUsOrErr) { - GpuArchs.insert(Args.MakeArgString(GPU)); - } - } else { - ArchStr = getCanonicalOffloadArch(ArchStr); - if (ArchStr.empty()) { - Error = true; - } else if (A->getOption().matches(options::OPT_offload_arch_EQ)) - GpuArchs.insert(ArchStr); - else if (A->getOption().matches(options::OPT_no_offload_arch_EQ)) - GpuArchs.erase(ArchStr); - else - llvm_unreachable("Unexpected option."); - } + for (auto Arch : + C.getDriver().getOffloadArchs(C, C.getArgs(), Kind, *I.second)) + GpuArchs.insert(Arch); } } - auto &&ConflictingArchs = getConflictOffloadArchCombination(GpuArchs); - if (ConflictingArchs) { - C.getDriver().Diag(clang::diag::err_drv_bad_offload_arch_combo) - << ConflictingArchs->first << ConflictingArchs->second; - C.setContainsError(); - return true; - } - - // Collect list of GPUs remaining in the set. for (auto Arch : GpuArchs) GpuArchList.push_back(Arch.data()); - // Default to sm_20 which is the lowest common denominator for - // supported GPUs. sm_20 code should work correctly, if - // suboptimally, on all newer GPUs. - if (GpuArchList.empty()) { - if (ToolChains.front()->getTriple().isSPIRV()) { - if (ToolChains.front()->getTriple().getVendor() == llvm::Triple::AMD) - GpuArchList.push_back(OffloadArch::AMDGCNSPIRV); - else - GpuArchList.push_back(OffloadArch::Generic); - } else { - GpuArchList.push_back(DefaultOffloadArch); - } - } + CompileHostOnly = C.getDriver().offloadHostOnly(); + EmitLLVM = Args.getLastArg(options::OPT_emit_llvm); + EmitAsm = Args.getLastArg(options::OPT_S); - return Error; + return false; } }; @@ -3538,15 +3462,6 @@ class OffloadingActionBuilder final { DefaultOffloadArch = OffloadArch::CudaDefault; } - StringRef getCanonicalOffloadArch(StringRef ArchStr) override { - OffloadArch Arch = StringToOffloadArch(ArchStr); - if (Arch == OffloadArch::UNKNOWN || !IsNVIDIAOffloadArch(Arch)) { - C.getDriver().Diag(clang::diag::err_drv_cuda_bad_gpu_arch) << ArchStr; - return StringRef(); - } - return OffloadArchToString(Arch); - } - std::optional<std::pair<llvm::StringRef, llvm::StringRef>> getConflictOffloadArchCombination( const std::set<StringRef> &GpuArchs) override { @@ -3705,24 +3620,6 @@ class OffloadingActionBuilder final { bool canUseBundlerUnbundler() const override { return true; } - StringRef getCanonicalOffloadArch(StringRef IdStr) override { - llvm::StringMap<bool> Features; - // getHIPOffloadTargetTriple() is known to return valid value as it has - // been called successfully in the CreateOffloadingDeviceToolChains(). - auto T = - (IdStr == "amdgcnspirv") - ? llvm::Triple("spirv64-amd-amdhsa") - : *getHIPOffloadTargetTriple(C.getDriver(), C.getInputArgs()); - auto ArchStr = parseTargetID(T, IdStr, &Features); - if (!ArchStr) { - C.getDriver().Diag(clang::diag::err_drv_bad_target_id) << IdStr; - C.setContainsError(); - return StringRef(); - } - auto CanId = getCanonicalTargetID(*ArchStr, Features); - return Args.MakeArgStringRef(CanId); - }; - std::optional<std::pair<llvm::StringRef, llvm::StringRef>> getConflictOffloadArchCombination( const std::set<StringRef> &GpuArchs) override { @@ -4715,23 +4612,20 @@ void Driver::BuildActions(Compilation &C, DerivedArgList &Args, static StringRef getCanonicalArchString(Compilation &C, const llvm::opt::DerivedArgList &Args, StringRef ArchStr, - const llvm::Triple &Triple, - bool SpecificToolchain) { + const llvm::Triple &Triple) { // Lookup the CUDA / HIP architecture string. Only report an error if we were // expecting the triple to be only NVPTX / AMDGPU. OffloadArch Arch = StringToOffloadArch(getProcessorFromTargetID(Triple, ArchStr)); if (Triple.isNVPTX() && (Arch == OffloadArch::UNKNOWN || !IsNVIDIAOffloadArch(Arch))) { - if (SpecificToolchain) - C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch) - << "CUDA" << ArchStr; + C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch) + << "CUDA" << ArchStr; return StringRef(); } else if (Triple.isAMDGPU() && (Arch == OffloadArch::UNKNOWN || !IsAMDOffloadArch(Arch))) { - if (SpecificToolchain) - C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch) - << "HIP" << ArchStr; + C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch) + << "HIP" << ArchStr; return StringRef(); } @@ -4767,11 +4661,7 @@ getConflictOffloadArchCombination(const llvm::DenseSet<StringRef> &Archs, llvm::SmallVector<StringRef> Driver::getOffloadArchs(Compilation &C, const llvm::opt::DerivedArgList &Args, - Action::OffloadKind Kind, const ToolChain *TC, - bool SpecificToolchain) const { - if (!TC) - TC = &C.getDefaultToolChain(); - + Action::OffloadKind Kind, const ToolChain &TC) const { // --offload and --offload-arch options are mutually exclusive. if (Args.hasArgNoClaim(options::OPT_offload_EQ) && Args.hasArgNoClaim(options::OPT_offload_arch_EQ, @@ -4784,48 +4674,44 @@ Driver::getOffloadArchs(Compilation &C, const llvm::opt::DerivedArgList &Args, } llvm::DenseSet<StringRef> Archs; - for (auto *Arg : C.getArgsForToolChain(TC, /*BoundArch=*/"", Kind)) { + for (auto *Arg : C.getArgsForToolChain(&TC, /*BoundArch=*/"", Kind)) { // Add or remove the seen architectures in order of appearance. If an // invalid architecture is given we simply exit. if (Arg->getOption().matches(options::OPT_offload_arch_EQ)) { for (StringRef Arch : Arg->getValues()) { if (Arch == "native" || Arch.empty()) { - auto GPUsOrErr = TC->getSystemGPUArchs(Args); + auto GPUsOrErr = TC.getSystemGPUArchs(Args); if (!GPUsOrErr) { - if (!SpecificToolchain) - llvm::consumeError(GPUsOrErr.takeError()); - else - TC->getDriver().Diag(diag::err_drv_undetermined_gpu_arch) - << llvm::Triple::getArchTypeName(TC->getArch()) - << llvm::toString(GPUsOrErr.takeError()) << "--offload-arch"; + TC.getDriver().Diag(diag::err_drv_undetermined_gpu_arch) + << llvm::Triple::getArchTypeName(TC.getArch()) + << llvm::toString(GPUsOrErr.takeError()) << "--offload-arch"; continue; } for (auto ArchStr : *GPUsOrErr) { - StringRef CanonicalStr = - getCanonicalArchString(C, Args, Args.MakeArgString(ArchStr), - TC->getTriple(), SpecificToolchain); + StringRef CanonicalStr = getCanonicalArchString( + C, Args, Args.MakeArgString(ArchStr), TC.getTriple()); if (!CanonicalStr.empty()) Archs.insert(CanonicalStr); - else if (SpecificToolchain) + else return llvm::SmallVector<StringRef>(); } } else { - StringRef CanonicalStr = getCanonicalArchString( - C, Args, Arch, TC->getTriple(), SpecificToolchain); + StringRef CanonicalStr = + getCanonicalArchString(C, Args, Arch, TC.getTriple()); if (!CanonicalStr.empty()) Archs.insert(CanonicalStr); - else if (SpecificToolchain) + else return llvm::SmallVector<StringRef>(); } } } else if (Arg->getOption().matches(options::OPT_no_offload_arch_EQ)) { - for (StringRef Arch : llvm::split(Arg->getValue(), ",")) { + for (StringRef Arch : Arg->getValues()) { if (Arch == "all") { Archs.clear(); } else { - StringRef ArchStr = getCanonicalArchString( - C, Args, Arch, TC->getTriple(), SpecificToolchain); + StringRef ArchStr = + getCanonicalArchString(C, Args, Arch, TC.getTriple()); Archs.erase(ArchStr); } } @@ -4833,28 +4719,30 @@ Driver::getOffloadArchs(Compilation &C, const llvm::opt::DerivedArgList &Args, } if (auto ConflictingArchs = - getConflictOffloadArchCombination(Archs, TC->getTriple())) + getConflictOffloadArchCombination(Archs, TC.getTriple())) C.getDriver().Diag(clang::diag::err_drv_bad_offload_arch_combo) << ConflictingArchs->first << ConflictingArchs->second; - // Skip filling defaults if we're just querying what is availible. - if (SpecificToolchain && Archs.empty()) { + // Fill in the default architectures if not provided explicitly. + if (Archs.empty()) { if (Kind == Action::OFK_Cuda) { Archs.insert(OffloadArchToString(OffloadArch::CudaDefault)); } else if (Kind == Action::OFK_HIP) { - Archs.insert(OffloadArchToString(OffloadArch::HIPDefault)); + Archs.insert(OffloadArchToString(TC.getTriple().isSPIRV() + ? OffloadArch::Generic + : OffloadArch::HIPDefault)); } else if (Kind == Action::OFK_SYCL) { Archs.insert(StringRef()); } else if (Kind == Action::OFK_OpenMP) { // Accept legacy `-march` device arguments for OpenMP. - if (auto *Arg = C.getArgsForToolChain(TC, /*BoundArch=*/"", Kind) + if (auto *Arg = C.getArgsForToolChain(&TC, /*BoundArch=*/"", Kind) .getLastArg(options::OPT_march_EQ)) { Archs.insert(Arg->getValue()); } else { - auto ArchsOrErr = TC->getSystemGPUArchs(Args); + auto ArchsOrErr = TC.getSystemGPUArchs(Args); if (!ArchsOrErr) { - TC->getDriver().Diag(diag::err_drv_undetermined_gpu_arch) - << llvm::Triple::getArchTypeName(TC->getArch()) + TC.getDriver().Diag(diag::err_drv_undetermined_gpu_arch) + << llvm::Triple::getArchTypeName(TC.getArch()) << llvm::toString(ArchsOrErr.takeError()) << "--offload-arch"; } else if (!ArchsOrErr->empty()) { for (auto Arch : *ArchsOrErr) @@ -4934,7 +4822,7 @@ Action *Driver::BuildOffloadingActions(Compilation &C, // Get the product of all bound architectures and toolchains. SmallVector<std::pair<const ToolChain *, StringRef>> TCAndArchs; for (const ToolChain *TC : ToolChains) { - for (StringRef Arch : OffloadArchs.lookup(TC)) { + for (StringRef Arch : getOffloadArchs(C, C.getArgs(), Kind, *TC)) { TCAndArchs.push_back(std::make_pair(TC, Arch)); DeviceActions.push_back( C.MakeAction<InputAction>(*InputArg, InputType, CUID)); @@ -4966,7 +4854,7 @@ Action *Driver::BuildOffloadingActions(Compilation &C, if (Kind == Action::OFK_SYCL && Phase == phases::Assemble) continue; - auto TCAndArch = TCAndArchs.begin(); + auto *TCAndArch = TCAndArchs.begin(); for (Action *&A : DeviceActions) { if (A->getType() == types::TY_Nothing) continue; @@ -4998,7 +4886,13 @@ Action *Driver::BuildOffloadingActions(Compilation &C, // Compiling HIP in device-only non-RDC mode requires linking each action // individually. for (Action *&A : DeviceActions) { - if ((A->getType() != types::TY_Object && + // Special handling for the HIP SPIR-V toolchain because it doesn't use + // the SPIR-V backend yet doesn't report the output as an object. + bool IsAMDGCNSPIRV = A->getOffloadingToolChain() && + A->getOffloadingToolChain()->getTriple().getOS() == + llvm::Triple::OSType::AMDHSA && + A->getOffloadingToolChain()->getTriple().isSPIRV(); + if ((A->getType() != types::TY_Object && !IsAMDGCNSPIRV && A->getType() != types::TY_LTO_BC) || !HIPNoRDC || !offloadDeviceOnly()) continue; @@ -5006,7 +4900,7 @@ Action *Driver::BuildOffloadingActions(Compilation &C, A = C.MakeAction<LinkJobAction>(LinkerInput, types::TY_Image); } - auto TCAndArch = TCAndArchs.begin(); + auto *TCAndArch = TCAndArchs.begin(); for (Action *A : DeviceActions) { DDeps.add(*A, *TCAndArch->first, TCAndArch->second.data(), Kind); OffloadAction::DeviceDependences DDep; @@ -5054,8 +4948,9 @@ Action *Driver::BuildOffloadingActions(Compilation &C, // fatbinary for each translation unit, linking each input individually. Action *FatbinAction = C.MakeAction<LinkJobAction>(OffloadActions, types::TY_HIP_FATBIN); - DDep.add(*FatbinAction, *C.getSingleOffloadToolChain<Action::OFK_HIP>(), - nullptr, Action::OFK_HIP); + DDep.add(*FatbinAction, + *C.getOffloadToolChains<Action::OFK_HIP>().first->second, nullptr, + Action::OFK_HIP); } else { // Package all the offloading actions into a single output that can be // embedded in the host and linked. @@ -5210,7 +5105,10 @@ Action *Driver::ConstructPhaseAction( false) || (Args.hasFlag(options::OPT_offload_new_driver, options::OPT_no_offload_new_driver, false) && - !offloadDeviceOnly())) || + (!offloadDeviceOnly() || + (Input->getOffloadingToolChain() && + TargetDeviceOffloadKind == Action::OFK_HIP && + Input->getOffloadingToolChain()->getTriple().isSPIRV())))) || TargetDeviceOffloadKind == Action::OFK_OpenMP))) { types::ID Output = Args.hasArg(options::OPT_S) && diff --git a/clang/lib/Driver/SanitizerArgs.cpp b/clang/lib/Driver/SanitizerArgs.cpp index 21e4cff..98793a5 100644 --- a/clang/lib/Driver/SanitizerArgs.cpp +++ b/clang/lib/Driver/SanitizerArgs.cpp @@ -1382,6 +1382,12 @@ void SanitizerArgs::addArgs(const ToolChain &TC, const llvm::opt::ArgList &Args, CmdArgs.push_back(Args.MakeArgString("-fsanitize-annotate-debug-info=" + toString(AnnotateDebugInfo))); + if (const Arg *A = + Args.getLastArg(options::OPT_fsanitize_debug_trap_reasons, + options::OPT_fno_sanitize_debug_trap_reasons)) { + CmdArgs.push_back(Args.MakeArgString(A->getAsString(Args))); + } + addSpecialCaseListOpt(Args, CmdArgs, "-fsanitize-ignorelist=", UserIgnorelistFiles); addSpecialCaseListOpt(Args, CmdArgs, diff --git a/clang/lib/Driver/ToolChain.cpp b/clang/lib/Driver/ToolChain.cpp index 481f575..1d7dad0 100644 --- a/clang/lib/Driver/ToolChain.cpp +++ b/clang/lib/Driver/ToolChain.cpp @@ -104,44 +104,6 @@ ToolChain::ToolChain(const Driver &D, const llvm::Triple &T, addIfExists(getFilePaths(), Path); } -llvm::Expected<std::unique_ptr<llvm::MemoryBuffer>> -ToolChain::executeToolChainProgram(StringRef Executable) const { - llvm::SmallString<64> OutputFile; - llvm::sys::fs::createTemporaryFile("toolchain-program", "txt", OutputFile, - llvm::sys::fs::OF_Text); - llvm::FileRemover OutputRemover(OutputFile.c_str()); - std::optional<llvm::StringRef> Redirects[] = { - {""}, - OutputFile.str(), - {""}, - }; - - std::string ErrorMessage; - int SecondsToWait = 60; - if (std::optional<std::string> Str = - llvm::sys::Process::GetEnv("CLANG_TOOLCHAIN_PROGRAM_TIMEOUT")) { - if (!llvm::to_integer(*Str, SecondsToWait)) - return llvm::createStringError(std::error_code(), - "CLANG_TOOLCHAIN_PROGRAM_TIMEOUT expected " - "an integer, got '" + - *Str + "'"); - SecondsToWait = std::max(SecondsToWait, 0); // infinite - } - if (llvm::sys::ExecuteAndWait(Executable, {Executable}, {}, Redirects, - SecondsToWait, - /*MemoryLimit=*/0, &ErrorMessage)) - return llvm::createStringError(std::error_code(), - Executable + ": " + ErrorMessage); - - llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> OutputBuf = - llvm::MemoryBuffer::getFile(OutputFile.c_str()); - if (!OutputBuf) - return llvm::createStringError(OutputBuf.getError(), - "Failed to read stdout of " + Executable + - ": " + OutputBuf.getError().message()); - return std::move(*OutputBuf); -} - void ToolChain::setTripleEnvironment(llvm::Triple::EnvironmentType Env) { Triple.setEnvironment(Env); if (EffectiveTriple != llvm::Triple()) @@ -255,6 +217,18 @@ static void getAArch64MultilibFlags(const Driver &D, Result.push_back(ABIArg->getAsString(Args)); } + if (const Arg *A = Args.getLastArg(options::OPT_O_Group); + A && A->getOption().matches(options::OPT_O)) { + switch (A->getValue()[0]) { + case 's': + Result.push_back("-Os"); + break; + case 'z': + Result.push_back("-Oz"); + break; + } + } + processMultilibCustomFlags(Result, Args); } @@ -332,6 +306,19 @@ static void getARMMultilibFlags(const Driver &D, const llvm::Triple &Triple, if (Endian->getOption().matches(options::OPT_mbig_endian)) Result.push_back(Endian->getAsString(Args)); } + + if (const Arg *A = Args.getLastArg(options::OPT_O_Group); + A && A->getOption().matches(options::OPT_O)) { + switch (A->getValue()[0]) { + case 's': + Result.push_back("-Os"); + break; + case 'z': + Result.push_back("-Oz"); + break; + } + } + processMultilibCustomFlags(Result, Args); } @@ -1100,7 +1087,7 @@ std::string ToolChain::GetLinkerPath(bool *LinkerIsLLD) const { // Get -fuse-ld= first to prevent -Wunused-command-line-argument. -fuse-ld= is // considered as the linker flavor, e.g. "bfd", "gold", or "lld". const Arg* A = Args.getLastArg(options::OPT_fuse_ld_EQ); - StringRef UseLinker = A ? A->getValue() : CLANG_DEFAULT_LINKER; + StringRef UseLinker = A ? A->getValue() : getDriver().getPreferredLinker(); // --ld-path= takes precedence over -fuse-ld= and specifies the executable // name. -B, COMPILER_PATH and PATH and consulted if the value does not @@ -1644,7 +1631,8 @@ void ToolChain::addSYCLIncludeArgs(const ArgList &DriverArgs, ArgStringList &CC1Args) const {} llvm::SmallVector<ToolChain::BitCodeLibraryInfo, 12> -ToolChain::getDeviceLibs(const ArgList &DriverArgs) const { +ToolChain::getDeviceLibs(const ArgList &DriverArgs, + const Action::OffloadKind DeviceOffloadingKind) const { return {}; } diff --git a/clang/lib/Driver/ToolChains/AMDGPU.cpp b/clang/lib/Driver/ToolChains/AMDGPU.cpp index 7fc34f4..0781683 100644 --- a/clang/lib/Driver/ToolChains/AMDGPU.cpp +++ b/clang/lib/Driver/ToolChains/AMDGPU.cpp @@ -31,6 +31,68 @@ using namespace clang::driver::toolchains; using namespace clang; using namespace llvm::opt; +RocmInstallationDetector::CommonBitcodeLibsPreferences:: + CommonBitcodeLibsPreferences(const Driver &D, + const llvm::opt::ArgList &DriverArgs, + StringRef GPUArch, + const Action::OffloadKind DeviceOffloadingKind, + const bool NeedsASanRT) + : ABIVer(DeviceLibABIVersion::fromCodeObjectVersion( + tools::getAMDGPUCodeObjectVersion(D, DriverArgs))) { + const auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch); + const unsigned ArchAttr = llvm::AMDGPU::getArchAttrAMDGCN(Kind); + + IsOpenMP = DeviceOffloadingKind == Action::OFK_OpenMP; + + const bool HasWave32 = (ArchAttr & llvm::AMDGPU::FEATURE_WAVE32); + Wave64 = + !HasWave32 || DriverArgs.hasFlag(options::OPT_mwavefrontsize64, + options::OPT_mno_wavefrontsize64, false); + + const bool IsKnownOffloading = DeviceOffloadingKind == Action::OFK_OpenMP || + DeviceOffloadingKind == Action::OFK_HIP; + + // Default to enabling f32 denormals on subtargets where fma is fast with + // denormals + const bool DefaultDAZ = + (Kind == llvm::AMDGPU::GK_NONE) + ? false + : !((ArchAttr & llvm::AMDGPU::FEATURE_FAST_FMA_F32) && + (ArchAttr & llvm::AMDGPU::FEATURE_FAST_DENORMAL_F32)); + // TODO: There are way too many flags that change this. Do we need to + // check them all? + DAZ = IsKnownOffloading + ? DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero, + options::OPT_fno_gpu_flush_denormals_to_zero, + DefaultDAZ) + : DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) || DefaultDAZ; + + FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only) || + DriverArgs.hasFlag(options::OPT_ffinite_math_only, + options::OPT_fno_finite_math_only, false); + + UnsafeMathOpt = + DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations) || + DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations, + options::OPT_fno_unsafe_math_optimizations, false); + + FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math) || + DriverArgs.hasFlag(options::OPT_ffast_math, + options::OPT_fno_fast_math, false); + + const bool DefaultSqrt = IsKnownOffloading ? true : false; + CorrectSqrt = + DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt) || + DriverArgs.hasFlag( + options::OPT_fhip_fp32_correctly_rounded_divide_sqrt, + options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt, DefaultSqrt); + // GPU Sanitizer currently only supports ASan and is enabled through host + // ASan. + GPUSan = (DriverArgs.hasFlag(options::OPT_fgpu_sanitize, + options::OPT_fno_gpu_sanitize, true) && + NeedsASanRT); +} + void RocmInstallationDetector::scanLibDevicePath(llvm::StringRef Path) { assert(!Path.empty()); @@ -841,7 +903,7 @@ AMDGPUToolChain::getSystemGPUArchs(const ArgList &Args) const { else Program = GetProgramPath("amdgpu-arch"); - auto StdoutOrErr = executeToolChainProgram(Program); + auto StdoutOrErr = getDriver().executeProgram({Program}); if (!StdoutOrErr) return StdoutOrErr.takeError(); @@ -884,33 +946,14 @@ void ROCMToolChain::addClangTargetOptions( ABIVer)) return; - bool Wave64 = isWave64(DriverArgs, Kind); - // TODO: There are way too many flags that change this. Do we need to check - // them all? - bool DAZ = DriverArgs.hasArg(options::OPT_cl_denorms_are_zero) || - getDefaultDenormsAreZeroForTarget(Kind); - bool FiniteOnly = DriverArgs.hasArg(options::OPT_cl_finite_math_only); - - bool UnsafeMathOpt = - DriverArgs.hasArg(options::OPT_cl_unsafe_math_optimizations); - bool FastRelaxedMath = DriverArgs.hasArg(options::OPT_cl_fast_relaxed_math); - bool CorrectSqrt = - DriverArgs.hasArg(options::OPT_cl_fp32_correctly_rounded_divide_sqrt); - - // GPU Sanitizer currently only supports ASan and is enabled through host - // ASan. - bool GPUSan = DriverArgs.hasFlag(options::OPT_fgpu_sanitize, - options::OPT_fno_gpu_sanitize, true) && - getSanitizerArgs(DriverArgs).needsAsanRt(); - // Add the OpenCL specific bitcode library. llvm::SmallVector<BitCodeLibraryInfo, 12> BCLibs; BCLibs.emplace_back(RocmInstallation->getOpenCLPath().str()); // Add the generic set of libraries. BCLibs.append(RocmInstallation->getCommonBitcodeLibs( - DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt, - FastRelaxedMath, CorrectSqrt, ABIVer, GPUSan, false)); + DriverArgs, LibDeviceFile, GpuArch, DeviceOffloadingKind, + getSanitizerArgs(DriverArgs).needsAsanRt())); for (auto [BCFile, Internalize] : BCLibs) { if (Internalize) @@ -947,35 +990,37 @@ bool RocmInstallationDetector::checkCommonBitcodeLibs( llvm::SmallVector<ToolChain::BitCodeLibraryInfo, 12> RocmInstallationDetector::getCommonBitcodeLibs( - const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, bool Wave64, - bool DAZ, bool FiniteOnly, bool UnsafeMathOpt, bool FastRelaxedMath, - bool CorrectSqrt, DeviceLibABIVersion ABIVer, bool GPUSan, - bool isOpenMP) const { + const llvm::opt::ArgList &DriverArgs, StringRef LibDeviceFile, + StringRef GPUArch, const Action::OffloadKind DeviceOffloadingKind, + const bool NeedsASanRT) const { llvm::SmallVector<ToolChain::BitCodeLibraryInfo, 12> BCLibs; + CommonBitcodeLibsPreferences Pref{D, DriverArgs, GPUArch, + DeviceOffloadingKind, NeedsASanRT}; + auto AddBCLib = [&](ToolChain::BitCodeLibraryInfo BCLib, bool Internalize = true) { BCLib.ShouldInternalize = Internalize; BCLibs.emplace_back(BCLib); }; auto AddSanBCLibs = [&]() { - if (GPUSan) + if (Pref.GPUSan) AddBCLib(getAsanRTLPath(), false); }; AddSanBCLibs(); AddBCLib(getOCMLPath()); - if (!isOpenMP) + if (!Pref.IsOpenMP) AddBCLib(getOCKLPath()); - else if (GPUSan && isOpenMP) + else if (Pref.GPUSan && Pref.IsOpenMP) AddBCLib(getOCKLPath(), false); - AddBCLib(getDenormalsAreZeroPath(DAZ)); - AddBCLib(getUnsafeMathPath(UnsafeMathOpt || FastRelaxedMath)); - AddBCLib(getFiniteOnlyPath(FiniteOnly || FastRelaxedMath)); - AddBCLib(getCorrectlyRoundedSqrtPath(CorrectSqrt)); - AddBCLib(getWavefrontSize64Path(Wave64)); + AddBCLib(getDenormalsAreZeroPath(Pref.DAZ)); + AddBCLib(getUnsafeMathPath(Pref.UnsafeMathOpt || Pref.FastRelaxedMath)); + AddBCLib(getFiniteOnlyPath(Pref.FiniteOnly || Pref.FastRelaxedMath)); + AddBCLib(getCorrectlyRoundedSqrtPath(Pref.CorrectSqrt)); + AddBCLib(getWavefrontSize64Path(Pref.Wave64)); AddBCLib(LibDeviceFile); - auto ABIVerPath = getABIVersionPath(ABIVer); + auto ABIVerPath = getABIVersionPath(Pref.ABIVer); if (!ABIVerPath.empty()) AddBCLib(ABIVerPath); @@ -983,9 +1028,9 @@ RocmInstallationDetector::getCommonBitcodeLibs( } llvm::SmallVector<ToolChain::BitCodeLibraryInfo, 12> -ROCMToolChain::getCommonDeviceLibNames(const llvm::opt::ArgList &DriverArgs, - const std::string &GPUArch, - bool isOpenMP) const { +ROCMToolChain::getCommonDeviceLibNames( + const llvm::opt::ArgList &DriverArgs, const std::string &GPUArch, + Action::OffloadKind DeviceOffloadingKind) const { auto Kind = llvm::AMDGPU::parseArchAMDGCN(GPUArch); const StringRef CanonArch = llvm::AMDGPU::getArchNameAMDGCN(Kind); @@ -996,33 +1041,9 @@ ROCMToolChain::getCommonDeviceLibNames(const llvm::opt::ArgList &DriverArgs, ABIVer)) return {}; - // If --hip-device-lib is not set, add the default bitcode libraries. - // TODO: There are way too many flags that change this. Do we need to check - // them all? - bool DAZ = DriverArgs.hasFlag(options::OPT_fgpu_flush_denormals_to_zero, - options::OPT_fno_gpu_flush_denormals_to_zero, - getDefaultDenormsAreZeroForTarget(Kind)); - bool FiniteOnly = DriverArgs.hasFlag( - options::OPT_ffinite_math_only, options::OPT_fno_finite_math_only, false); - bool UnsafeMathOpt = - DriverArgs.hasFlag(options::OPT_funsafe_math_optimizations, - options::OPT_fno_unsafe_math_optimizations, false); - bool FastRelaxedMath = DriverArgs.hasFlag(options::OPT_ffast_math, - options::OPT_fno_fast_math, false); - bool CorrectSqrt = DriverArgs.hasFlag( - options::OPT_fhip_fp32_correctly_rounded_divide_sqrt, - options::OPT_fno_hip_fp32_correctly_rounded_divide_sqrt, true); - bool Wave64 = isWave64(DriverArgs, Kind); - - // GPU Sanitizer currently only supports ASan and is enabled through host - // ASan. - bool GPUSan = DriverArgs.hasFlag(options::OPT_fgpu_sanitize, - options::OPT_fno_gpu_sanitize, true) && - getSanitizerArgs(DriverArgs).needsAsanRt(); - return RocmInstallation->getCommonBitcodeLibs( - DriverArgs, LibDeviceFile, Wave64, DAZ, FiniteOnly, UnsafeMathOpt, - FastRelaxedMath, CorrectSqrt, ABIVer, GPUSan, isOpenMP); + DriverArgs, LibDeviceFile, GPUArch, DeviceOffloadingKind, + getSanitizerArgs(DriverArgs).needsAsanRt()); } bool AMDGPUToolChain::shouldSkipSanitizeOption( diff --git a/clang/lib/Driver/ToolChains/AMDGPU.h b/clang/lib/Driver/ToolChains/AMDGPU.h index 08bd4fa..513c77d 100644 --- a/clang/lib/Driver/ToolChains/AMDGPU.h +++ b/clang/lib/Driver/ToolChains/AMDGPU.h @@ -147,7 +147,7 @@ public: llvm::SmallVector<BitCodeLibraryInfo, 12> getCommonDeviceLibNames(const llvm::opt::ArgList &DriverArgs, const std::string &GPUArch, - bool isOpenMP = false) const; + Action::OffloadKind DeviceOffloadingKind) const; SanitizerMask getSupportedSanitizers() const override { return SanitizerKind::Address; diff --git a/clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp b/clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp index 7ffa3f0..2b41d54 100644 --- a/clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp +++ b/clang/lib/Driver/ToolChains/AMDGPUOpenMP.cpp @@ -44,7 +44,7 @@ void AMDGPUOpenMPToolChain::addClangTargetOptions( true)) return; - for (auto BCFile : getDeviceLibs(DriverArgs)) { + for (auto BCFile : getDeviceLibs(DriverArgs, DeviceOffloadingKind)) { CC1Args.push_back(BCFile.ShouldInternalize ? "-mlink-builtin-bitcode" : "-mlink-bitcode-file"); CC1Args.push_back(DriverArgs.MakeArgString(BCFile.Path)); @@ -132,7 +132,9 @@ AMDGPUOpenMPToolChain::computeMSVCVersion(const Driver *D, } llvm::SmallVector<ToolChain::BitCodeLibraryInfo, 12> -AMDGPUOpenMPToolChain::getDeviceLibs(const llvm::opt::ArgList &Args) const { +AMDGPUOpenMPToolChain::getDeviceLibs( + const llvm::opt::ArgList &Args, + const Action::OffloadKind DeviceOffloadingKind) const { if (!Args.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib, true)) return {}; @@ -140,8 +142,8 @@ AMDGPUOpenMPToolChain::getDeviceLibs(const llvm::opt::ArgList &Args) const { getTriple(), Args.getLastArgValue(options::OPT_march_EQ)); SmallVector<BitCodeLibraryInfo, 12> BCLibs; - for (auto BCLib : getCommonDeviceLibNames(Args, GpuArch.str(), - /*IsOpenMP=*/true)) + for (auto BCLib : + getCommonDeviceLibNames(Args, GpuArch.str(), DeviceOffloadingKind)) BCLibs.emplace_back(BCLib); return BCLibs; diff --git a/clang/lib/Driver/ToolChains/AMDGPUOpenMP.h b/clang/lib/Driver/ToolChains/AMDGPUOpenMP.h index 0536c9f..cbafdf5 100644 --- a/clang/lib/Driver/ToolChains/AMDGPUOpenMP.h +++ b/clang/lib/Driver/ToolChains/AMDGPUOpenMP.h @@ -58,7 +58,8 @@ public: const llvm::opt::ArgList &Args) const override; llvm::SmallVector<BitCodeLibraryInfo, 12> - getDeviceLibs(const llvm::opt::ArgList &Args) const override; + getDeviceLibs(const llvm::opt::ArgList &Args, + const Action::OffloadKind DeviceOffloadKind) const override; const ToolChain &HostTC; }; diff --git a/clang/lib/Driver/ToolChains/Arch/Sparc.cpp b/clang/lib/Driver/ToolChains/Arch/Sparc.cpp index 504f110..94a94f1 100644 --- a/clang/lib/Driver/ToolChains/Arch/Sparc.cpp +++ b/clang/lib/Driver/ToolChains/Arch/Sparc.cpp @@ -23,7 +23,9 @@ const char *sparc::getSparcAsmModeForCPU(StringRef Name, if (Triple.getArch() == llvm::Triple::sparcv9) { const char *DefV9CPU; - if (Triple.isOSLinux() || Triple.isOSFreeBSD() || Triple.isOSOpenBSD()) + if (Triple.isOSSolaris()) + DefV9CPU = "-Av9b"; + else if (Triple.isOSLinux() || Triple.isOSFreeBSD() || Triple.isOSOpenBSD()) DefV9CPU = "-Av9a"; else DefV9CPU = "-Av9"; @@ -35,6 +37,13 @@ const char *sparc::getSparcAsmModeForCPU(StringRef Name, .Case("niagara4", "-Av9d") .Default(DefV9CPU); } else { + const char *DefV8CPU; + + if (Triple.isOSSolaris()) + DefV8CPU = "-Av8plus"; + else + DefV8CPU = "-Av8"; + return llvm::StringSwitch<const char *>(Name) .Case("v8", "-Av8") .Case("supersparc", "-Av8") @@ -70,7 +79,7 @@ const char *sparc::getSparcAsmModeForCPU(StringRef Name, .Case("gr712rc", "-Aleon") .Case("leon4", "-Aleon") .Case("gr740", "-Aleon") - .Default("-Av8"); + .Default(DefV8CPU); } } @@ -157,6 +166,9 @@ void sparc::getSparcTargetFeatures(const Driver &D, const llvm::Triple &Triple, bool IsSparcV9ATarget = (Triple.getArch() == llvm::Triple::sparcv9) && (Triple.isOSLinux() || Triple.isOSFreeBSD() || Triple.isOSOpenBSD()); + bool IsSparcV9BTarget = Triple.isOSSolaris(); + bool IsSparcV8PlusTarget = + Triple.getArch() == llvm::Triple::sparc && Triple.isOSSolaris(); if (Arg *A = Args.getLastArg(options::OPT_mvis, options::OPT_mno_vis)) { if (A->getOption().matches(options::OPT_mvis)) Features.push_back("+vis"); @@ -171,6 +183,8 @@ void sparc::getSparcTargetFeatures(const Driver &D, const llvm::Triple &Triple, Features.push_back("+vis2"); else Features.push_back("-vis2"); + } else if (IsSparcV9BTarget) { + Features.push_back("+vis2"); } if (Arg *A = Args.getLastArg(options::OPT_mvis3, options::OPT_mno_vis3)) { @@ -191,6 +205,8 @@ void sparc::getSparcTargetFeatures(const Driver &D, const llvm::Triple &Triple, if (Arg *A = Args.getLastArg(options::OPT_mv8plus, options::OPT_mno_v8plus)) { if (A->getOption().matches(options::OPT_mv8plus)) Features.push_back("+v8plus"); + } else if (IsSparcV8PlusTarget) { + Features.push_back("+v8plus"); } if (Args.hasArg(options::OPT_ffixed_g1)) diff --git a/clang/lib/Driver/ToolChains/BareMetal.cpp b/clang/lib/Driver/ToolChains/BareMetal.cpp index e670696..497f333 100644 --- a/clang/lib/Driver/ToolChains/BareMetal.cpp +++ b/clang/lib/Driver/ToolChains/BareMetal.cpp @@ -694,9 +694,6 @@ void baremetal::Linker::ConstructJob(Compilation &C, const JobAction &JA, NeedCRTs) CmdArgs.push_back(Args.MakeArgString(TC.GetFilePath(CRTEnd))); - if (TC.getTriple().isRISCV()) - CmdArgs.push_back("-X"); - // The R_ARM_TARGET2 relocation must be treated as R_ARM_REL32 on arm*-*-elf // and arm*-*-eabi (the default is R_ARM_GOT_PREL, used on arm*-*-linux and // arm*-*-*bsd). diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 1fc7002..9d882db 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -97,32 +97,15 @@ forAllAssociatedToolChains(Compilation &C, const JobAction &JA, // Apply Work on all the offloading tool chains associated with the current // action. - if (JA.isHostOffloading(Action::OFK_Cuda)) - Work(*C.getSingleOffloadToolChain<Action::OFK_Cuda>()); - else if (JA.isDeviceOffloading(Action::OFK_Cuda)) - Work(*C.getSingleOffloadToolChain<Action::OFK_Host>()); - else if (JA.isHostOffloading(Action::OFK_HIP)) - Work(*C.getSingleOffloadToolChain<Action::OFK_HIP>()); - else if (JA.isDeviceOffloading(Action::OFK_HIP)) - Work(*C.getSingleOffloadToolChain<Action::OFK_Host>()); - - if (JA.isHostOffloading(Action::OFK_OpenMP)) { - auto TCs = C.getOffloadToolChains<Action::OFK_OpenMP>(); - for (auto II = TCs.first, IE = TCs.second; II != IE; ++II) - Work(*II->second); - } else if (JA.isDeviceOffloading(Action::OFK_OpenMP)) - Work(*C.getSingleOffloadToolChain<Action::OFK_Host>()); - - if (JA.isHostOffloading(Action::OFK_SYCL)) { - auto TCs = C.getOffloadToolChains<Action::OFK_SYCL>(); - for (auto II = TCs.first, IE = TCs.second; II != IE; ++II) - Work(*II->second); - } else if (JA.isDeviceOffloading(Action::OFK_SYCL)) - Work(*C.getSingleOffloadToolChain<Action::OFK_Host>()); - - // - // TODO: Add support for other offloading programming models here. - // + for (Action::OffloadKind Kind : {Action::OFK_Cuda, Action::OFK_OpenMP, + Action::OFK_HIP, Action::OFK_SYCL}) { + if (JA.isHostOffloading(Kind)) { + auto TCs = C.getOffloadToolChains(Kind); + for (auto II = TCs.first, IE = TCs.second; II != IE; ++II) + Work(*II->second); + } else if (JA.isDeviceOffloading(Kind)) + Work(*C.getSingleOffloadToolChain<Action::OFK_Host>()); + } } static bool @@ -3898,17 +3881,17 @@ static bool RenderModulesOptions(Compilation &C, const Driver &D, const ArgList &Args, const InputInfo &Input, const InputInfo &Output, bool HaveStd20, ArgStringList &CmdArgs) { - bool IsCXX = types::isCXX(Input.getType()); - bool HaveStdCXXModules = IsCXX && HaveStd20; + const bool IsCXX = types::isCXX(Input.getType()); + const bool HaveStdCXXModules = IsCXX && HaveStd20; bool HaveModules = HaveStdCXXModules; // -fmodules enables the use of precompiled modules (off by default). // Users can pass -fno-cxx-modules to turn off modules support for // C++/Objective-C++ programs. + const bool AllowedInCXX = Args.hasFlag(options::OPT_fcxx_modules, + options::OPT_fno_cxx_modules, true); bool HaveClangModules = false; if (Args.hasFlag(options::OPT_fmodules, options::OPT_fno_modules, false)) { - bool AllowedInCXX = Args.hasFlag(options::OPT_fcxx_modules, - options::OPT_fno_cxx_modules, true); if (AllowedInCXX || !IsCXX) { CmdArgs.push_back("-fmodules"); HaveClangModules = true; @@ -3917,6 +3900,9 @@ static bool RenderModulesOptions(Compilation &C, const Driver &D, HaveModules |= HaveClangModules; + if (HaveModules && !AllowedInCXX) + CmdArgs.push_back("-fno-cxx-modules"); + // -fmodule-maps enables implicit reading of module map files. By default, // this is enabled if we are using Clang's flavor of precompiled modules. if (Args.hasFlag(options::OPT_fimplicit_module_maps, @@ -4985,8 +4971,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, else { // Host-side compilation. NormalizedTriple = - (IsCuda ? C.getSingleOffloadToolChain<Action::OFK_Cuda>() - : C.getSingleOffloadToolChain<Action::OFK_HIP>()) + (IsCuda ? C.getOffloadToolChains(Action::OFK_Cuda).first->second + : C.getOffloadToolChains(Action::OFK_HIP).first->second) ->getTriple() .normalize(); if (IsCuda) { diff --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp index 2373d94..7d803be 100644 --- a/clang/lib/Driver/ToolChains/Cuda.cpp +++ b/clang/lib/Driver/ToolChains/Cuda.cpp @@ -815,7 +815,7 @@ NVPTXToolChain::getSystemGPUArchs(const ArgList &Args) const { else Program = GetProgramPath("nvptx-arch"); - auto StdoutOrErr = executeToolChainProgram(Program); + auto StdoutOrErr = getDriver().executeProgram({Program}); if (!StdoutOrErr) return StdoutOrErr.takeError(); diff --git a/clang/lib/Driver/ToolChains/HIPAMD.cpp b/clang/lib/Driver/ToolChains/HIPAMD.cpp index 5fe0f85..b4c6da0 100644 --- a/clang/lib/Driver/ToolChains/HIPAMD.cpp +++ b/clang/lib/Driver/ToolChains/HIPAMD.cpp @@ -264,7 +264,7 @@ void HIPAMDToolChain::addClangTargetOptions( return; // No DeviceLibs for SPIR-V. } - for (auto BCFile : getDeviceLibs(DriverArgs)) { + for (auto BCFile : getDeviceLibs(DriverArgs, DeviceOffloadingKind)) { CC1Args.push_back(BCFile.ShouldInternalize ? "-mlink-builtin-bitcode" : "-mlink-bitcode-file"); CC1Args.push_back(DriverArgs.MakeArgString(BCFile.Path)); @@ -355,7 +355,8 @@ VersionTuple HIPAMDToolChain::computeMSVCVersion(const Driver *D, } llvm::SmallVector<ToolChain::BitCodeLibraryInfo, 12> -HIPAMDToolChain::getDeviceLibs(const llvm::opt::ArgList &DriverArgs) const { +HIPAMDToolChain::getDeviceLibs(const llvm::opt::ArgList &DriverArgs, + Action::OffloadKind DeviceOffloadingKind) const { llvm::SmallVector<BitCodeLibraryInfo, 12> BCLibs; if (!DriverArgs.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib, true) || @@ -397,7 +398,8 @@ HIPAMDToolChain::getDeviceLibs(const llvm::opt::ArgList &DriverArgs) const { assert(!GpuArch.empty() && "Must have an explicit GPU arch."); // Add common device libraries like ocml etc. - for (auto N : getCommonDeviceLibNames(DriverArgs, GpuArch.str())) + for (auto N : getCommonDeviceLibNames(DriverArgs, GpuArch.str(), + DeviceOffloadingKind)) BCLibs.emplace_back(N); // Add instrument lib. diff --git a/clang/lib/Driver/ToolChains/HIPAMD.h b/clang/lib/Driver/ToolChains/HIPAMD.h index 3630b11..bcc3ebb 100644 --- a/clang/lib/Driver/ToolChains/HIPAMD.h +++ b/clang/lib/Driver/ToolChains/HIPAMD.h @@ -80,7 +80,8 @@ public: void AddHIPIncludeArgs(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args) const override; llvm::SmallVector<BitCodeLibraryInfo, 12> - getDeviceLibs(const llvm::opt::ArgList &Args) const override; + getDeviceLibs(const llvm::opt::ArgList &Args, + Action::OffloadKind DeviceOffloadKind) const override; SanitizerMask getSupportedSanitizers() const override; diff --git a/clang/lib/Driver/ToolChains/HIPSPV.cpp b/clang/lib/Driver/ToolChains/HIPSPV.cpp index 53649ca..643a67f 100644 --- a/clang/lib/Driver/ToolChains/HIPSPV.cpp +++ b/clang/lib/Driver/ToolChains/HIPSPV.cpp @@ -149,7 +149,8 @@ void HIPSPVToolChain::addClangTargetOptions( CC1Args.append( {"-fvisibility=hidden", "-fapply-global-visibility-to-externs"}); - for (const BitCodeLibraryInfo &BCFile : getDeviceLibs(DriverArgs)) + for (const BitCodeLibraryInfo &BCFile : + getDeviceLibs(DriverArgs, DeviceOffloadingKind)) CC1Args.append( {"-mlink-builtin-bitcode", DriverArgs.MakeArgString(BCFile.Path)}); } @@ -200,7 +201,9 @@ void HIPSPVToolChain::AddHIPIncludeArgs(const ArgList &DriverArgs, } llvm::SmallVector<ToolChain::BitCodeLibraryInfo, 12> -HIPSPVToolChain::getDeviceLibs(const llvm::opt::ArgList &DriverArgs) const { +HIPSPVToolChain::getDeviceLibs( + const llvm::opt::ArgList &DriverArgs, + const Action::OffloadKind DeviceOffloadingKind) const { llvm::SmallVector<ToolChain::BitCodeLibraryInfo, 12> BCLibs; if (!DriverArgs.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib, true)) diff --git a/clang/lib/Driver/ToolChains/HIPSPV.h b/clang/lib/Driver/ToolChains/HIPSPV.h index ecd82e7..caf6924 100644 --- a/clang/lib/Driver/ToolChains/HIPSPV.h +++ b/clang/lib/Driver/ToolChains/HIPSPV.h @@ -69,7 +69,8 @@ public: void AddHIPIncludeArgs(const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args) const override; llvm::SmallVector<BitCodeLibraryInfo, 12> - getDeviceLibs(const llvm::opt::ArgList &Args) const override; + getDeviceLibs(const llvm::opt::ArgList &Args, + const Action::OffloadKind DeviceOffloadKind) const override; SanitizerMask getSupportedSanitizers() const override; diff --git a/clang/lib/Driver/ToolChains/MSVC.cpp b/clang/lib/Driver/ToolChains/MSVC.cpp index 7d31eea..bb469ff 100644 --- a/clang/lib/Driver/ToolChains/MSVC.cpp +++ b/clang/lib/Driver/ToolChains/MSVC.cpp @@ -279,8 +279,8 @@ void visualstudio::Linker::ConstructJob(Compilation &C, const JobAction &JA, AddRunTimeLibs(TC, TC.getDriver(), CmdArgs, Args); } - StringRef Linker = - Args.getLastArgValue(options::OPT_fuse_ld_EQ, CLANG_DEFAULT_LINKER); + StringRef Linker = Args.getLastArgValue(options::OPT_fuse_ld_EQ, + TC.getDriver().getPreferredLinker()); if (Linker.empty()) Linker = "link"; // We need to translate 'lld' into 'lld-link'. diff --git a/clang/lib/Driver/ToolChains/MinGW.cpp b/clang/lib/Driver/ToolChains/MinGW.cpp index b2e36ae..6abd0c0 100644 --- a/clang/lib/Driver/ToolChains/MinGW.cpp +++ b/clang/lib/Driver/ToolChains/MinGW.cpp @@ -548,7 +548,7 @@ toolchains::MinGW::MinGW(const Driver &D, const llvm::Triple &Triple, getFilePaths().push_back(Base + "lib"); NativeLLVMSupport = - Args.getLastArgValue(options::OPT_fuse_ld_EQ, CLANG_DEFAULT_LINKER) + Args.getLastArgValue(options::OPT_fuse_ld_EQ, D.getPreferredLinker()) .equals_insensitive("lld"); } diff --git a/clang/lib/Driver/ToolChains/OpenBSD.cpp b/clang/lib/Driver/ToolChains/OpenBSD.cpp index 79b1b69..8f58918 100644 --- a/clang/lib/Driver/ToolChains/OpenBSD.cpp +++ b/clang/lib/Driver/ToolChains/OpenBSD.cpp @@ -161,7 +161,7 @@ void openbsd::Linker::ConstructJob(Compilation &C, const JobAction &JA, if (Nopie || Profiling) CmdArgs.push_back("-nopie"); - if (Triple.isRISCV64()) { + if (Triple.isLoongArch64() || Triple.isRISCV64()) { CmdArgs.push_back("-X"); if (Args.hasArg(options::OPT_mno_relax)) CmdArgs.push_back("--no-relax"); diff --git a/clang/lib/Driver/ToolChains/ROCm.h b/clang/lib/Driver/ToolChains/ROCm.h index 2a09da01..ebd5443 100644 --- a/clang/lib/Driver/ToolChains/ROCm.h +++ b/clang/lib/Driver/ToolChains/ROCm.h @@ -11,6 +11,7 @@ #include "clang/Basic/Cuda.h" #include "clang/Basic/LLVM.h" +#include "clang/Driver/CommonArgs.h" #include "clang/Driver/Driver.h" #include "clang/Driver/Options.h" #include "clang/Driver/SanitizerArgs.h" @@ -18,6 +19,7 @@ #include "llvm/ADT/StringMap.h" #include "llvm/Option/ArgList.h" #include "llvm/Support/VersionTuple.h" +#include "llvm/TargetParser/TargetParser.h" #include "llvm/TargetParser/Triple.h" namespace clang { @@ -77,6 +79,24 @@ private: SPACKReleaseStr(SPACKReleaseStr.str()) {} }; + struct CommonBitcodeLibsPreferences { + CommonBitcodeLibsPreferences(const Driver &D, + const llvm::opt::ArgList &DriverArgs, + StringRef GPUArch, + const Action::OffloadKind DeviceOffloadingKind, + const bool NeedsASanRT); + + DeviceLibABIVersion ABIVer; + bool IsOpenMP; + bool Wave64; + bool DAZ; + bool FiniteOnly; + bool UnsafeMathOpt; + bool FastRelaxedMath; + bool CorrectSqrt; + bool GPUSan; + }; + const Driver &D; bool HasHIPRuntime = false; bool HasDeviceLibrary = false; @@ -175,11 +195,11 @@ public: /// Get file paths of default bitcode libraries common to AMDGPU based /// toolchains. - 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, - bool GPUSan, bool isOpenMP) const; + llvm::SmallVector<ToolChain::BitCodeLibraryInfo, 12> + getCommonBitcodeLibs(const llvm::opt::ArgList &DriverArgs, + StringRef LibDeviceFile, StringRef GPUArch, + const Action::OffloadKind DeviceOffloadingKind, + const bool NeedsASanRT) const; /// Check file paths of default bitcode libraries common to AMDGPU based /// toolchains. \returns false if there are invalid or missing files. bool checkCommonBitcodeLibs(StringRef GPUArch, StringRef LibDeviceFile, diff --git a/clang/lib/Driver/ToolChains/Solaris.cpp b/clang/lib/Driver/ToolChains/Solaris.cpp index a3574e1..02aa598 100644 --- a/clang/lib/Driver/ToolChains/Solaris.cpp +++ b/clang/lib/Driver/ToolChains/Solaris.cpp @@ -39,7 +39,7 @@ void solaris::Assembler::ConstructJob(Compilation &C, const JobAction &JA, bool solaris::isLinkerGnuLd(const ToolChain &TC, const ArgList &Args) { // Only used if targetting Solaris. const Arg *A = Args.getLastArg(options::OPT_fuse_ld_EQ); - StringRef UseLinker = A ? A->getValue() : CLANG_DEFAULT_LINKER; + StringRef UseLinker = A ? A->getValue() : TC.getDriver().getPreferredLinker(); return UseLinker == "bfd" || UseLinker == "gld"; } @@ -52,7 +52,7 @@ static bool getPIE(const ArgList &Args, const ToolChain &TC) { TC.isPIEDefault(Args)); } -// FIXME: Need to handle CLANG_DEFAULT_LINKER here? +// FIXME: Need to handle PreferredLinker here? std::string solaris::Linker::getLinkerPath(const ArgList &Args) const { const ToolChain &ToolChain = getToolChain(); if (const Arg *A = Args.getLastArg(options::OPT_fuse_ld_EQ)) { @@ -345,7 +345,7 @@ SanitizerMask Solaris::getSupportedSanitizers() const { const char *Solaris::getDefaultLinker() const { // FIXME: Only handle Solaris ld and GNU ld here. - return llvm::StringSwitch<const char *>(CLANG_DEFAULT_LINKER) + return llvm::StringSwitch<const char *>(getDriver().getPreferredLinker()) .Cases("bfd", "gld", "/usr/gnu/bin/ld") .Default("/usr/bin/ld"); } diff --git a/clang/lib/Driver/ToolChains/UEFI.cpp b/clang/lib/Driver/ToolChains/UEFI.cpp index ac6668e..2b41173 100644 --- a/clang/lib/Driver/ToolChains/UEFI.cpp +++ b/clang/lib/Driver/ToolChains/UEFI.cpp @@ -83,8 +83,8 @@ void tools::uefi::Linker::ConstructJob(Compilation &C, const JobAction &JA, // This should ideally be handled by ToolChain::GetLinkerPath but we need // to special case some linker paths. In the case of lld, we need to // translate 'lld' into 'lld-link'. - StringRef Linker = - Args.getLastArgValue(options::OPT_fuse_ld_EQ, CLANG_DEFAULT_LINKER); + StringRef Linker = Args.getLastArgValue(options::OPT_fuse_ld_EQ, + TC.getDriver().getPreferredLinker()); if (Linker.empty() || Linker == "lld") Linker = "lld-link"; |