summaryrefslogtreecommitdiff
path: root/external/apache2/llvm/dist/clang/lib/CodeGen/CodeGenModule.cpp
diff options
context:
space:
mode:
authorjoerg <joerg@NetBSD.org>2021-05-30 01:25:15 +0000
committerjoerg <joerg@NetBSD.org>2021-05-30 01:25:15 +0000
commit13fbcb42fff2606c3ce926a7c4bb8989aa18758c (patch)
treedcace6c89967106796d44b1a7c5e2269a20967e2 /external/apache2/llvm/dist/clang/lib/CodeGen/CodeGenModule.cpp
parent13338e33506ade63440b4ba74479551bdc21c0f6 (diff)
Import clang 249b40b558955afe5ac2b549edcf2d7f859c8cc9.
Diffstat (limited to 'external/apache2/llvm/dist/clang/lib/CodeGen/CodeGenModule.cpp')
-rw-r--r--external/apache2/llvm/dist/clang/lib/CodeGen/CodeGenModule.cpp1576
1 files changed, 1032 insertions, 544 deletions
diff --git a/external/apache2/llvm/dist/clang/lib/CodeGen/CodeGenModule.cpp b/external/apache2/llvm/dist/clang/lib/CodeGen/CodeGenModule.cpp
index 480a33f2728..744ccf0fe59 100644
--- a/external/apache2/llvm/dist/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/external/apache2/llvm/dist/clang/lib/CodeGen/CodeGenModule.cpp
@@ -19,6 +19,7 @@
#include "CGObjCRuntime.h"
#include "CGOpenCLRuntime.h"
#include "CGOpenMPRuntime.h"
+#include "CGOpenMPRuntimeAMDGCN.h"
#include "CGOpenMPRuntimeNVPTX.h"
#include "CodeGenFunction.h"
#include "CodeGenPGO.h"
@@ -38,6 +39,7 @@
#include "clang/Basic/CharInfo.h"
#include "clang/Basic/CodeGenOptions.h"
#include "clang/Basic/Diagnostic.h"
+#include "clang/Basic/FileManager.h"
#include "clang/Basic/Module.h"
#include "clang/Basic/SourceManager.h"
#include "clang/Basic/TargetInfo.h"
@@ -47,6 +49,7 @@
#include "llvm/ADT/StringSwitch.h"
#include "llvm/ADT/Triple.h"
#include "llvm/Analysis/TargetLibraryInfo.h"
+#include "llvm/Frontend/OpenMP/OMPIRBuilder.h"
#include "llvm/IR/CallingConv.h"
#include "llvm/IR/DataLayout.h"
#include "llvm/IR/Intrinsics.h"
@@ -55,6 +58,7 @@
#include "llvm/IR/ProfileSummary.h"
#include "llvm/ProfileData/InstrProfReader.h"
#include "llvm/Support/CodeGen.h"
+#include "llvm/Support/CommandLine.h"
#include "llvm/Support/ConvertUTF.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/MD5.h"
@@ -71,15 +75,17 @@ static llvm::cl::opt<bool> LimitedCoverage(
static const char AnnotationSection[] = "llvm.metadata";
static CGCXXABI *createCXXABI(CodeGenModule &CGM) {
- switch (CGM.getTarget().getCXXABI().getKind()) {
+ switch (CGM.getContext().getCXXABIKind()) {
+ case TargetCXXABI::AppleARM64:
+ case TargetCXXABI::Fuchsia:
case TargetCXXABI::GenericAArch64:
case TargetCXXABI::GenericARM:
case TargetCXXABI::iOS:
- case TargetCXXABI::iOS64:
case TargetCXXABI::WatchOS:
case TargetCXXABI::GenericMIPS:
case TargetCXXABI::GenericItanium:
case TargetCXXABI::WebAssembly:
+ case TargetCXXABI::XL:
return CreateItaniumCXXABI(CGM);
case TargetCXXABI::Microsoft:
return CreateMicrosoftCXXABI(CGM);
@@ -107,6 +113,7 @@ CodeGenModule::CodeGenModule(ASTContext &C, const HeaderSearchOptions &HSO,
Int32Ty = llvm::Type::getInt32Ty(LLVMContext);
Int64Ty = llvm::Type::getInt64Ty(LLVMContext);
HalfTy = llvm::Type::getHalfTy(LLVMContext);
+ BFloatTy = llvm::Type::getBFloatTy(LLVMContext);
FloatTy = llvm::Type::getFloatTy(LLVMContext);
DoubleTy = llvm::Type::getDoubleTy(LLVMContext);
PointerWidthInBits = C.getTargetInfo().getPointerWidth(0);
@@ -116,6 +123,8 @@ CodeGenModule::CodeGenModule(ASTContext &C, const HeaderSearchOptions &HSO,
C.toCharUnitsFromBits(C.getTargetInfo().getMaxPointerWidth()).getQuantity();
IntAlignInBytes =
C.toCharUnitsFromBits(C.getTargetInfo().getIntAlign()).getQuantity();
+ CharTy =
+ llvm::IntegerType::get(LLVMContext, C.getTargetInfo().getCharWidth());
IntTy = llvm::IntegerType::get(LLVMContext, C.getTargetInfo().getIntWidth());
IntPtrTy = llvm::IntegerType::get(LLVMContext,
C.getTargetInfo().getMaxPointerWidth());
@@ -171,6 +180,34 @@ CodeGenModule::CodeGenModule(ASTContext &C, const HeaderSearchOptions &HSO,
// CoverageMappingModuleGen object.
if (CodeGenOpts.CoverageMapping)
CoverageMapping.reset(new CoverageMappingModuleGen(*this, *CoverageInfo));
+
+ // Generate the module name hash here if needed.
+ if (CodeGenOpts.UniqueInternalLinkageNames &&
+ !getModule().getSourceFileName().empty()) {
+ std::string Path = getModule().getSourceFileName();
+ // Check if a path substitution is needed from the MacroPrefixMap.
+ for (const auto &Entry : PPO.MacroPrefixMap)
+ if (Path.rfind(Entry.first, 0) != std::string::npos) {
+ Path = Entry.second + Path.substr(Entry.first.size());
+ break;
+ }
+ llvm::MD5 Md5;
+ Md5.update(Path);
+ llvm::MD5::MD5Result R;
+ Md5.final(R);
+ SmallString<32> Str;
+ llvm::MD5::stringifyResult(R, Str);
+ // Convert MD5hash to Decimal. Demangler suffixes can either contain
+ // numbers or characters but not both.
+ llvm::APInt IntHash(128, Str.str(), 16);
+ // Prepend "__uniq" before the hash for tools like profilers to understand
+ // that this symbol is of internal linkage type. The "__uniq" is the
+ // pre-determined prefix that is used to tell tools that this symbol was
+ // created with -funique-internal-linakge-symbols and the tools can strip or
+ // keep the prefix as needed.
+ ModuleNameHash = (Twine(".__uniq.") +
+ Twine(IntHash.toString(/* Radix = */ 10, /* Signed = */false))).str();
+ }
}
CodeGenModule::~CodeGenModule() {}
@@ -209,6 +246,11 @@ void CodeGenModule::createOpenMPRuntime() {
"OpenMP NVPTX is only prepared to deal with device code.");
OpenMPRuntime.reset(new CGOpenMPRuntimeNVPTX(*this));
break;
+ case llvm::Triple::amdgcn:
+ assert(getLangOpts().OpenMPIsDevice &&
+ "OpenMP AMDGCN is only prepared to deal with device code.");
+ OpenMPRuntime.reset(new CGOpenMPRuntimeAMDGCN(*this));
+ break;
default:
if (LangOpts.OpenMPSimd)
OpenMPRuntime.reset(new CGOpenMPSIMDRuntime(*this));
@@ -360,7 +402,7 @@ void CodeGenModule::checkAliases() {
for (const GlobalDecl &GD : Aliases) {
StringRef MangledName = getMangledName(GD);
llvm::GlobalValue *Entry = GetGlobalValue(MangledName);
- auto *Alias = dyn_cast<llvm::GlobalIndirectSymbol>(Entry);
+ auto *Alias = cast<llvm::GlobalIndirectSymbol>(Entry);
Alias->replaceAllUsesWith(llvm::UndefValue::get(Alias->getType()));
Alias->eraseFromParent();
}
@@ -389,6 +431,48 @@ void InstrProfStats::reportDiagnostics(DiagnosticsEngine &Diags,
}
}
+static void setVisibilityFromDLLStorageClass(const clang::LangOptions &LO,
+ llvm::Module &M) {
+ if (!LO.VisibilityFromDLLStorageClass)
+ return;
+
+ llvm::GlobalValue::VisibilityTypes DLLExportVisibility =
+ CodeGenModule::GetLLVMVisibility(LO.getDLLExportVisibility());
+ llvm::GlobalValue::VisibilityTypes NoDLLStorageClassVisibility =
+ CodeGenModule::GetLLVMVisibility(LO.getNoDLLStorageClassVisibility());
+ llvm::GlobalValue::VisibilityTypes ExternDeclDLLImportVisibility =
+ CodeGenModule::GetLLVMVisibility(LO.getExternDeclDLLImportVisibility());
+ llvm::GlobalValue::VisibilityTypes ExternDeclNoDLLStorageClassVisibility =
+ CodeGenModule::GetLLVMVisibility(
+ LO.getExternDeclNoDLLStorageClassVisibility());
+
+ for (llvm::GlobalValue &GV : M.global_values()) {
+ if (GV.hasAppendingLinkage() || GV.hasLocalLinkage())
+ continue;
+
+ // Reset DSO locality before setting the visibility. This removes
+ // any effects that visibility options and annotations may have
+ // had on the DSO locality. Setting the visibility will implicitly set
+ // appropriate globals to DSO Local; however, this will be pessimistic
+ // w.r.t. to the normal compiler IRGen.
+ GV.setDSOLocal(false);
+
+ if (GV.isDeclarationForLinker()) {
+ GV.setVisibility(GV.getDLLStorageClass() ==
+ llvm::GlobalValue::DLLImportStorageClass
+ ? ExternDeclDLLImportVisibility
+ : ExternDeclNoDLLStorageClassVisibility);
+ } else {
+ GV.setVisibility(GV.getDLLStorageClass() ==
+ llvm::GlobalValue::DLLExportStorageClass
+ ? DLLExportVisibility
+ : NoDLLStorageClassVisibility);
+ }
+
+ GV.setDLLStorageClass(llvm::GlobalValue::DefaultStorageClass);
+ }
+}
+
void CodeGenModule::Release() {
EmitDeferred();
EmitVTablesOpportunistically();
@@ -397,16 +481,14 @@ void CodeGenModule::Release() {
checkAliases();
emitMultiVersionFunctions();
EmitCXXGlobalInitFunc();
- EmitCXXGlobalDtorFunc();
+ EmitCXXGlobalCleanUpFunc();
registerGlobalDtorsWithAtExit();
EmitCXXThreadLocalInitFunc();
if (ObjCRuntime)
if (llvm::Function *ObjCInitFunction = ObjCRuntime->ModuleInitFunction())
AddGlobalCtor(ObjCInitFunction);
- if (Context.getLangOpts().CUDA && !Context.getLangOpts().CUDAIsDevice &&
- CUDARuntime) {
- if (llvm::Function *CudaCtorFunction =
- CUDARuntime->makeModuleCtorFunction())
+ if (Context.getLangOpts().CUDA && CUDARuntime) {
+ if (llvm::Function *CudaCtorFunction = CUDARuntime->finalizeModule())
AddGlobalCtor(CudaCtorFunction);
}
if (OpenMPRuntime) {
@@ -429,6 +511,7 @@ void CodeGenModule::Release() {
EmitGlobalAnnotations();
EmitStaticExternCAliases();
EmitDeferredUnusedCoverageMappings();
+ CodeGenPGO(*this).setValueProfilingFlag(getModule());
if (CoverageMapping)
CoverageMapping->emit();
if (CodeGenOpts.SanitizeCfiCrossDso) {
@@ -436,6 +519,10 @@ void CodeGenModule::Release() {
CodeGenFunction(*this).EmitCfiCheckStub();
}
emitAtAvailableLinkGuard();
+ if (Context.getTargetInfo().getTriple().isWasm() &&
+ !Context.getTargetInfo().getTriple().isOSEmscripten()) {
+ EmitMainVoidAlias();
+ }
emitLLVMUsed();
if (SanStats)
SanStats->finish();
@@ -469,11 +556,17 @@ void CodeGenModule::Release() {
CodeGenOpts.NumRegisterParameters);
if (CodeGenOpts.DwarfVersion) {
- // We actually want the latest version when there are conflicts.
- // We can change from Warning to Latest if such mode is supported.
- getModule().addModuleFlag(llvm::Module::Warning, "Dwarf Version",
+ getModule().addModuleFlag(llvm::Module::Max, "Dwarf Version",
CodeGenOpts.DwarfVersion);
}
+
+ if (CodeGenOpts.Dwarf64)
+ getModule().addModuleFlag(llvm::Module::Max, "DWARF64", 1);
+
+ if (Context.getLangOpts().SemanticInterposition)
+ // Require various optimization to respect semantic interposition.
+ getModule().setSemanticInterposition(1);
+
if (CodeGenOpts.EmitCodeView) {
// Indicate that we want CodeView in the metadata.
getModule().addModuleFlag(llvm::Module::Warning, "CodeView", 1);
@@ -488,6 +581,10 @@ void CodeGenModule::Release() {
// Function ID tables for Control Flow Guard (cfguard=1).
getModule().addModuleFlag(llvm::Module::Warning, "cfguard", 1);
}
+ if (CodeGenOpts.EHContGuard) {
+ // Function ID tables for EH Continuation Guard.
+ getModule().addModuleFlag(llvm::Module::Warning, "ehcontguard", 1);
+ }
if (CodeGenOpts.OptimizationLevel > 0 && CodeGenOpts.StrictVTablePointers) {
// We don't support LTO with 2 with different StrictVTablePointers
// FIXME: we could support it by stripping all the information introduced
@@ -504,7 +601,7 @@ void CodeGenModule::Release() {
"StrictVTablePointersRequirement",
llvm::MDNode::get(VMContext, Ops));
}
- if (DebugInfo)
+ if (getModuleDebugInfo())
// We support a single version in the linked module. The LLVM
// parser will drop debug info with a different version number
// (and warn about it, too).
@@ -528,11 +625,26 @@ void CodeGenModule::Release() {
getModule().addModuleFlag(llvm::Module::Error, "min_enum_size", EnumWidth);
}
+ if (Arch == llvm::Triple::riscv32 || Arch == llvm::Triple::riscv64) {
+ StringRef ABIStr = Target.getABI();
+ llvm::LLVMContext &Ctx = TheModule.getContext();
+ getModule().addModuleFlag(llvm::Module::Error, "target-abi",
+ llvm::MDString::get(Ctx, ABIStr));
+ }
+
if (CodeGenOpts.SanitizeCfiCrossDso) {
// Indicate that we want cross-DSO control flow integrity checks.
getModule().addModuleFlag(llvm::Module::Override, "Cross-DSO CFI", 1);
}
+ if (CodeGenOpts.WholeProgramVTables) {
+ // Indicate whether VFE was enabled for this module, so that the
+ // vcall_visibility metadata added under whole program vtables is handled
+ // appropriately in the optimizer.
+ getModule().addModuleFlag(llvm::Module::Error, "Virtual Function Elim",
+ CodeGenOpts.VirtualFunctionElimination);
+ }
+
if (LangOpts.Sanitize.has(SanitizerKind::CFIICall)) {
getModule().addModuleFlag(llvm::Module::Override,
"CFI Canonical Jump Tables",
@@ -553,14 +665,42 @@ void CodeGenModule::Release() {
1);
}
+ if (Arch == llvm::Triple::aarch64 || Arch == llvm::Triple::aarch64_32 ||
+ Arch == llvm::Triple::aarch64_be) {
+ getModule().addModuleFlag(llvm::Module::Error,
+ "branch-target-enforcement",
+ LangOpts.BranchTargetEnforcement);
+
+ getModule().addModuleFlag(llvm::Module::Error, "sign-return-address",
+ LangOpts.hasSignReturnAddress());
+
+ getModule().addModuleFlag(llvm::Module::Error, "sign-return-address-all",
+ LangOpts.isSignReturnAddressScopeAll());
+
+ getModule().addModuleFlag(llvm::Module::Error,
+ "sign-return-address-with-bkey",
+ !LangOpts.isSignReturnAddressWithAKey());
+ }
+
+ if (!CodeGenOpts.MemoryProfileOutput.empty()) {
+ llvm::LLVMContext &Ctx = TheModule.getContext();
+ getModule().addModuleFlag(
+ llvm::Module::Error, "MemProfProfileFilename",
+ llvm::MDString::get(Ctx, CodeGenOpts.MemoryProfileOutput));
+ }
+
if (LangOpts.CUDAIsDevice && getTriple().isNVPTX()) {
// Indicate whether __nvvm_reflect should be configured to flush denormal
// floating point values to 0. (This corresponds to its "__CUDA_FTZ"
// property.)
getModule().addModuleFlag(llvm::Module::Override, "nvvm-reflect-ftz",
- CodeGenOpts.FlushDenorm ? 1 : 0);
+ CodeGenOpts.FP32DenormalMode.Output !=
+ llvm::DenormalMode::IEEE);
}
+ if (LangOpts.EHAsynch)
+ getModule().addModuleFlag(llvm::Module::Warning, "eh-asynch", 1);
+
// Emit OpenCL specific module metadata: OpenCL/SPIR version.
if (LangOpts.OpenCL) {
EmitOpenCLMetadata();
@@ -605,6 +745,20 @@ void CodeGenModule::Release() {
if (CodeGenOpts.NoPLT)
getModule().setRtLibUseGOT();
+ if (CodeGenOpts.UnwindTables)
+ getModule().setUwtable();
+
+ switch (CodeGenOpts.getFramePointer()) {
+ case CodeGenOptions::FramePointerKind::None:
+ // 0 ("none") is the default.
+ break;
+ case CodeGenOptions::FramePointerKind::NonLeaf:
+ getModule().setFramePointer(llvm::FramePointerKind::NonLeaf);
+ break;
+ case CodeGenOptions::FramePointerKind::All:
+ getModule().setFramePointer(llvm::FramePointerKind::All);
+ break;
+ }
SimplifyPersonality();
@@ -614,8 +768,8 @@ void CodeGenModule::Release() {
if (getCodeGenOpts().EmitGcovArcs || getCodeGenOpts().EmitGcovNotes)
EmitCoverageFile();
- if (DebugInfo)
- DebugInfo->finalize();
+ if (CGDebugInfo *DI = getModuleDebugInfo())
+ DI->finalize();
if (getCodeGenOpts().EmitVersionIdentMetadata)
EmitVersionIdentMetadata();
@@ -623,7 +777,24 @@ void CodeGenModule::Release() {
if (!getCodeGenOpts().RecordCommandLine.empty())
EmitCommandLineMetadata();
- EmitTargetMetadata();
+ if (!getCodeGenOpts().StackProtectorGuard.empty())
+ getModule().setStackProtectorGuard(getCodeGenOpts().StackProtectorGuard);
+ if (!getCodeGenOpts().StackProtectorGuardReg.empty())
+ getModule().setStackProtectorGuardReg(
+ getCodeGenOpts().StackProtectorGuardReg);
+ if (getCodeGenOpts().StackProtectorGuardOffset != INT_MAX)
+ getModule().setStackProtectorGuardOffset(
+ getCodeGenOpts().StackProtectorGuardOffset);
+
+ getTargetCodeGenInfo().emitTargetMetadata(*this, MangledDeclNames);
+
+ EmitBackendOptionsMetadata(getCodeGenOpts());
+
+ // Set visibility from DLL storage class
+ // We do this at the end of LLVM IR generation; after any operation
+ // that might affect the DLL storage class or the visibility, and
+ // before anything that might act on these.
+ setVisibilityFromDLLStorageClass(LangOpts, getModule());
}
void CodeGenModule::EmitOpenCLMetadata() {
@@ -643,6 +814,19 @@ void CodeGenModule::EmitOpenCLMetadata() {
OCLVerMD->addOperand(llvm::MDNode::get(Ctx, OCLVerElts));
}
+void CodeGenModule::EmitBackendOptionsMetadata(
+ const CodeGenOptions CodeGenOpts) {
+ switch (getTriple().getArch()) {
+ default:
+ break;
+ case llvm::Triple::riscv32:
+ case llvm::Triple::riscv64:
+ getModule().addModuleFlag(llvm::Module::Error, "SmallDataLimit",
+ CodeGenOpts.SmallDataLimit);
+ break;
+ }
+}
+
void CodeGenModule::UpdateCompletedType(const TagDecl *TD) {
// Make sure that this type is translated.
Types.UpdateCompletedType(TD);
@@ -662,6 +846,19 @@ llvm::MDNode *CodeGenModule::getTBAATypeInfo(QualType QTy) {
TBAAAccessInfo CodeGenModule::getTBAAAccessInfo(QualType AccessType) {
if (!TBAA)
return TBAAAccessInfo();
+ if (getLangOpts().CUDAIsDevice) {
+ // As CUDA builtin surface/texture types are replaced, skip generating TBAA
+ // access info.
+ if (AccessType->isCUDADeviceBuiltinSurfaceType()) {
+ if (getTargetCodeGenInfo().getCUDADeviceBuiltinSurfaceDeviceType() !=
+ nullptr)
+ return TBAAAccessInfo();
+ } else if (AccessType->isCUDADeviceBuiltinTextureType()) {
+ if (getTargetCodeGenInfo().getCUDADeviceBuiltinTextureDeviceType() !=
+ nullptr)
+ return TBAAAccessInfo();
+ }
+ }
return TBAA->getAccessInfo(AccessType);
}
@@ -816,8 +1013,17 @@ static bool shouldAssumeDSOLocal(const CodeGenModule &CGM,
const auto &CGOpts = CGM.getCodeGenOpts();
llvm::Reloc::Model RM = CGOpts.RelocationModel;
const auto &LOpts = CGM.getLangOpts();
- if (RM != llvm::Reloc::Static && !LOpts.PIE && !LOpts.OpenMPIsDevice)
- return false;
+ if (RM != llvm::Reloc::Static && !LOpts.PIE) {
+ // On ELF, if -fno-semantic-interposition is specified and the target
+ // supports local aliases, there will be neither CC1
+ // -fsemantic-interposition nor -fhalf-no-semantic-interposition. Set
+ // dso_local on the function if using a local alias is preferable (can avoid
+ // PLT indirection).
+ if (!(isa<llvm::Function>(GV) && GV->canBenefitFromLocalAlias()))
+ return false;
+ return !(CGM.getLangOpts().SemanticInterposition ||
+ CGM.getLangOpts().HalfNoSemanticInterposition);
+ }
// A definition cannot be preempted from an executable.
if (!GV->isDeclarationForLinker())
@@ -829,25 +1035,33 @@ static bool shouldAssumeDSOLocal(const CodeGenModule &CGM,
if (RM == llvm::Reloc::PIC_ && GV->hasExternalWeakLinkage())
return false;
- // PPC has no copy relocations and cannot use a plt entry as a symbol address.
- llvm::Triple::ArchType Arch = TT.getArch();
- if (Arch == llvm::Triple::ppc || Arch == llvm::Triple::ppc64 ||
- Arch == llvm::Triple::ppc64le)
+ // PowerPC64 prefers TOC indirection to avoid copy relocations.
+ if (TT.isPPC64())
return false;
- // If we can use copy relocations we can assume it is local.
- if (auto *Var = dyn_cast<llvm::GlobalVariable>(GV))
- if (!Var->isThreadLocal() &&
- (RM == llvm::Reloc::Static || CGOpts.PIECopyRelocations))
+ if (CGOpts.DirectAccessExternalData) {
+ // If -fdirect-access-external-data (default for -fno-pic), set dso_local
+ // for non-thread-local variables. If the symbol is not defined in the
+ // executable, a copy relocation will be needed at link time. dso_local is
+ // excluded for thread-local variables because they generally don't support
+ // copy relocations.
+ if (auto *Var = dyn_cast<llvm::GlobalVariable>(GV))
+ if (!Var->isThreadLocal())
+ return true;
+
+ // -fno-pic sets dso_local on a function declaration to allow direct
+ // accesses when taking its address (similar to a data symbol). If the
+ // function is not defined in the executable, a canonical PLT entry will be
+ // needed at link time. -fno-direct-access-external-data can avoid the
+ // canonical PLT entry. We don't generalize this condition to -fpie/-fpic as
+ // it could just cause trouble without providing perceptible benefits.
+ if (isa<llvm::Function>(GV) && !CGOpts.NoPLT && RM == llvm::Reloc::Static)
return true;
+ }
- // If we can use a plt entry as the symbol address we can assume it
- // is local.
- // FIXME: This should work for PIE, but the gold linker doesn't support it.
- if (isa<llvm::Function>(GV) && !CGOpts.NoPLT && RM == llvm::Reloc::Static)
- return true;
+ // If we can use copy relocations we can assume it is local.
- // Otherwise don't assue it is local.
+ // Otherwise don't assume it is local.
return false;
}
@@ -903,9 +1117,9 @@ static llvm::GlobalVariable::ThreadLocalMode GetLLVMTLSModel(StringRef S) {
.Case("local-exec", llvm::GlobalVariable::LocalExecTLSModel);
}
-static llvm::GlobalVariable::ThreadLocalMode GetLLVMTLSModel(
- CodeGenOptions::TLSModel M) {
- switch (M) {
+llvm::GlobalVariable::ThreadLocalMode
+CodeGenModule::GetDefaultLLVMTLSModel() const {
+ switch (CodeGenOpts.getDefaultTLSModel()) {
case CodeGenOptions::GeneralDynamicTLSModel:
return llvm::GlobalVariable::GeneralDynamicTLSModel;
case CodeGenOptions::LocalDynamicTLSModel:
@@ -922,7 +1136,7 @@ void CodeGenModule::setTLSMode(llvm::GlobalValue *GV, const VarDecl &D) const {
assert(D.getTLSKind() && "setting TLS mode on non-TLS var!");
llvm::GlobalValue::ThreadLocalMode TLM;
- TLM = GetLLVMTLSModel(CodeGenOpts.getDefaultTLSModel());
+ TLM = GetDefaultLLVMTLSModel();
// Override the TLS model if it is explicitly specified.
if (const TLSModelAttr *Attr = D.getAttr<TLSModelAttr>()) {
@@ -957,7 +1171,7 @@ static void AppendTargetMangling(const CodeGenModule &CGM,
Out << '.';
const TargetInfo &Target = CGM.getTarget();
- TargetAttr::ParsedTargetAttr Info =
+ ParsedTargetAttr Info =
Attr->parse([&Target](StringRef LHS, StringRef RHS) {
// Multiversioning doesn't allow "no-${feature}", so we can
// only have "+" prefixes here.
@@ -982,34 +1196,56 @@ static void AppendTargetMangling(const CodeGenModule &CGM,
}
}
-static std::string getMangledNameImpl(const CodeGenModule &CGM, GlobalDecl GD,
+// Returns true if GD is a function decl with internal linkage and
+// needs a unique suffix after the mangled name.
+static bool isUniqueInternalLinkageDecl(GlobalDecl GD,
+ CodeGenModule &CGM) {
+ const Decl *D = GD.getDecl();
+ return !CGM.getModuleNameHash().empty() && isa<FunctionDecl>(D) &&
+ (CGM.getFunctionLinkage(GD) == llvm::GlobalValue::InternalLinkage);
+}
+
+static std::string getMangledNameImpl(CodeGenModule &CGM, GlobalDecl GD,
const NamedDecl *ND,
bool OmitMultiVersionMangling = false) {
SmallString<256> Buffer;
llvm::raw_svector_ostream Out(Buffer);
MangleContext &MC = CGM.getCXXABI().getMangleContext();
- if (MC.shouldMangleDeclName(ND)) {
- llvm::raw_svector_ostream Out(Buffer);
- if (const auto *D = dyn_cast<CXXConstructorDecl>(ND))
- MC.mangleCXXCtor(D, GD.getCtorType(), Out);
- else if (const auto *D = dyn_cast<CXXDestructorDecl>(ND))
- MC.mangleCXXDtor(D, GD.getDtorType(), Out);
- else
- MC.mangleName(ND, Out);
- } else {
+ if (!CGM.getModuleNameHash().empty())
+ MC.needsUniqueInternalLinkageNames();
+ bool ShouldMangle = MC.shouldMangleDeclName(ND);
+ if (ShouldMangle)
+ MC.mangleName(GD.getWithDecl(ND), Out);
+ else {
IdentifierInfo *II = ND->getIdentifier();
assert(II && "Attempt to mangle unnamed decl.");
const auto *FD = dyn_cast<FunctionDecl>(ND);
if (FD &&
FD->getType()->castAs<FunctionType>()->getCallConv() == CC_X86RegCall) {
- llvm::raw_svector_ostream Out(Buffer);
Out << "__regcall3__" << II->getName();
+ } else if (FD && FD->hasAttr<CUDAGlobalAttr>() &&
+ GD.getKernelReferenceKind() == KernelReferenceKind::Stub) {
+ Out << "__device_stub__" << II->getName();
} else {
Out << II->getName();
}
}
+ // Check if the module name hash should be appended for internal linkage
+ // symbols. This should come before multi-version target suffixes are
+ // appended. This is to keep the name and module hash suffix of the
+ // internal linkage function together. The unique suffix should only be
+ // added when name mangling is done to make sure that the final name can
+ // be properly demangled. For example, for C functions without prototypes,
+ // name mangling is not done and the unique suffix should not be appeneded
+ // then.
+ if (ShouldMangle && isUniqueInternalLinkageDecl(GD, CGM)) {
+ assert(CGM.getCodeGenOpts().UniqueInternalLinkageNames &&
+ "Hash computed when not explicitly requested");
+ Out << CGM.getModuleNameHash();
+ }
+
if (const auto *FD = dyn_cast<FunctionDecl>(ND))
if (FD->isMultiVersion() && !OmitMultiVersionMangling) {
switch (FD->getMultiVersionKind()) {
@@ -1027,7 +1263,12 @@ static std::string getMangledNameImpl(const CodeGenModule &CGM, GlobalDecl GD,
}
}
- return Out.str();
+ // Make unique name for device side static file-scope variable for HIP.
+ if (CGM.getContext().shouldExternalizeStaticVar(ND) &&
+ CGM.getLangOpts().GPURelocatableDeviceCode &&
+ CGM.getLangOpts().CUDAIsDevice && !CGM.getLangOpts().CUID.empty())
+ CGM.printPostfixForExternalizedStaticVar(Out);
+ return std::string(Out.str());
}
void CodeGenModule::UpdateMultiVersionNames(GlobalDecl GD,
@@ -1084,19 +1325,40 @@ StringRef CodeGenModule::getMangledName(GlobalDecl GD) {
}
}
- auto FoundName = MangledDeclNames.find(CanonicalGD);
- if (FoundName != MangledDeclNames.end())
- return FoundName->second;
+ // In CUDA/HIP device compilation with -fgpu-rdc, the mangled name of a
+ // static device variable depends on whether the variable is referenced by
+ // a host or device host function. Therefore the mangled name cannot be
+ // cached.
+ if (!LangOpts.CUDAIsDevice ||
+ !getContext().mayExternalizeStaticVar(GD.getDecl())) {
+ auto FoundName = MangledDeclNames.find(CanonicalGD);
+ if (FoundName != MangledDeclNames.end())
+ return FoundName->second;
+ }
// Keep the first result in the case of a mangling collision.
const auto *ND = cast<NamedDecl>(GD.getDecl());
std::string MangledName = getMangledNameImpl(*this, GD, ND);
- // Adjust kernel stub mangling as we may need to be able to differentiate
- // them from the kernel itself (e.g., for HIP).
- if (auto *FD = dyn_cast<FunctionDecl>(GD.getDecl()))
- if (!getLangOpts().CUDAIsDevice && FD->hasAttr<CUDAGlobalAttr>())
- MangledName = getCUDARuntime().getDeviceStubName(MangledName);
+ // Ensure either we have different ABIs between host and device compilations,
+ // says host compilation following MSVC ABI but device compilation follows
+ // Itanium C++ ABI or, if they follow the same ABI, kernel names after
+ // mangling should be the same after name stubbing. The later checking is
+ // very important as the device kernel name being mangled in host-compilation
+ // is used to resolve the device binaries to be executed. Inconsistent naming
+ // result in undefined behavior. Even though we cannot check that naming
+ // directly between host- and device-compilations, the host- and
+ // device-mangling in host compilation could help catching certain ones.
+ assert(!isa<FunctionDecl>(ND) || !ND->hasAttr<CUDAGlobalAttr>() ||
+ getLangOpts().CUDAIsDevice ||
+ (getContext().getAuxTargetInfo() &&
+ (getContext().getAuxTargetInfo()->getCXXABI() !=
+ getContext().getTargetInfo().getCXXABI())) ||
+ getCUDARuntime().getDeviceSideName(ND) ==
+ getMangledNameImpl(
+ *this,
+ GD.getWithKernelReferenceKind(KernelReferenceKind::Kernel),
+ ND));
auto Result = Manglings.insert(std::make_pair(MangledName, GD));
return MangledDeclNames[CanonicalGD] = Result.first->first();
@@ -1137,8 +1399,10 @@ void CodeGenModule::AddGlobalCtor(llvm::Function *Ctor, int Priority,
/// AddGlobalDtor - Add a function to the list that will be called
/// when the module is unloaded.
-void CodeGenModule::AddGlobalDtor(llvm::Function *Dtor, int Priority) {
- if (CodeGenOpts.RegisterGlobalDtorsWithAtExit) {
+void CodeGenModule::AddGlobalDtor(llvm::Function *Dtor, int Priority,
+ bool IsDtorAttrFunc) {
+ if (CodeGenOpts.RegisterGlobalDtorsWithAtExit &&
+ (!getContext().getTargetInfo().getTriple().isOSAIX() || IsDtorAttrFunc)) {
DtorsUsingAtExit[Priority].push_back(Dtor);
return;
}
@@ -1215,10 +1479,11 @@ llvm::ConstantInt *CodeGenModule::CreateCrossDsoCfiTypeId(llvm::Metadata *MD) {
void CodeGenModule::SetLLVMFunctionAttributes(GlobalDecl GD,
const CGFunctionInfo &Info,
- llvm::Function *F) {
+ llvm::Function *F, bool IsThunk) {
unsigned CallingConv;
llvm::AttributeList PAL;
- ConstructAttributeList(F->getName(), Info, GD, PAL, CallingConv, false);
+ ConstructAttributeList(F->getName(), Info, GD, PAL, CallingConv,
+ /*AttrOnCallSite=*/false, IsThunk);
F->setAttributes(PAL);
F->setCallingConv(static_cast<llvm::CallingConv::ID>(CallingConv));
}
@@ -1251,10 +1516,18 @@ static void removeImageAccessQualifier(std::string& TyName) {
// (basically all single AS CPUs).
static unsigned ArgInfoAddressSpace(LangAS AS) {
switch (AS) {
- case LangAS::opencl_global: return 1;
- case LangAS::opencl_constant: return 2;
- case LangAS::opencl_local: return 3;
- case LangAS::opencl_generic: return 4; // Not in SPIR 2.0 specs.
+ case LangAS::opencl_global:
+ return 1;
+ case LangAS::opencl_constant:
+ return 2;
+ case LangAS::opencl_local:
+ return 3;
+ case LangAS::opencl_generic:
+ return 4; // Not in SPIR 2.0 specs.
+ case LangAS::opencl_global_device:
+ return 5;
+ case LangAS::opencl_global_host:
+ return 6;
default:
return 0; // Assume private.
}
@@ -1295,6 +1568,39 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
QualType ty = parm->getType();
std::string typeQuals;
+ // Get image and pipe access qualifier:
+ if (ty->isImageType() || ty->isPipeType()) {
+ const Decl *PDecl = parm;
+ if (auto *TD = dyn_cast<TypedefType>(ty))
+ PDecl = TD->getDecl();
+ const OpenCLAccessAttr *A = PDecl->getAttr<OpenCLAccessAttr>();
+ if (A && A->isWriteOnly())
+ accessQuals.push_back(llvm::MDString::get(VMContext, "write_only"));
+ else if (A && A->isReadWrite())
+ accessQuals.push_back(llvm::MDString::get(VMContext, "read_write"));
+ else
+ accessQuals.push_back(llvm::MDString::get(VMContext, "read_only"));
+ } else
+ accessQuals.push_back(llvm::MDString::get(VMContext, "none"));
+
+ // Get argument name.
+ argNames.push_back(llvm::MDString::get(VMContext, parm->getName()));
+
+ auto getTypeSpelling = [&](QualType Ty) {
+ auto typeName = Ty.getUnqualifiedType().getAsString(Policy);
+
+ if (Ty.isCanonical()) {
+ StringRef typeNameRef = typeName;
+ // Turn "unsigned type" to "utype"
+ if (typeNameRef.consume_front("unsigned "))
+ return std::string("u") + typeNameRef.str();
+ if (typeNameRef.consume_front("signed "))
+ return typeNameRef.str();
+ }
+
+ return typeName;
+ };
+
if (ty->isPointerType()) {
QualType pointeeTy = ty->getPointeeType();
@@ -1304,26 +1610,10 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
ArgInfoAddressSpace(pointeeTy.getAddressSpace()))));
// Get argument type name.
- std::string typeName =
- pointeeTy.getUnqualifiedType().getAsString(Policy) + "*";
-
- // Turn "unsigned type" to "utype"
- std::string::size_type pos = typeName.find("unsigned");
- if (pointeeTy.isCanonical() && pos != std::string::npos)
- typeName.erase(pos + 1, 8);
-
- argTypeNames.push_back(llvm::MDString::get(VMContext, typeName));
-
+ std::string typeName = getTypeSpelling(pointeeTy) + "*";
std::string baseTypeName =
- pointeeTy.getUnqualifiedType().getCanonicalType().getAsString(
- Policy) +
- "*";
-
- // Turn "unsigned type" to "utype"
- pos = baseTypeName.find("unsigned");
- if (pos != std::string::npos)
- baseTypeName.erase(pos + 1, 8);
-
+ getTypeSpelling(pointeeTy.getCanonicalType()) + "*";
+ argTypeNames.push_back(llvm::MDString::get(VMContext, typeName));
argBaseTypeNames.push_back(
llvm::MDString::get(VMContext, baseTypeName));
@@ -1345,30 +1635,9 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(AddrSpc)));
// Get argument type name.
- std::string typeName;
- if (isPipe)
- typeName = ty.getCanonicalType()
- ->getAs<PipeType>()
- ->getElementType()
- .getAsString(Policy);
- else
- typeName = ty.getUnqualifiedType().getAsString(Policy);
-
- // Turn "unsigned type" to "utype"
- std::string::size_type pos = typeName.find("unsigned");
- if (ty.isCanonical() && pos != std::string::npos)
- typeName.erase(pos + 1, 8);
-
- std::string baseTypeName;
- if (isPipe)
- baseTypeName = ty.getCanonicalType()
- ->getAs<PipeType>()
- ->getElementType()
- .getCanonicalType()
- .getAsString(Policy);
- else
- baseTypeName =
- ty.getUnqualifiedType().getCanonicalType().getAsString(Policy);
+ ty = isPipe ? ty->castAs<PipeType>()->getElementType() : ty;
+ std::string typeName = getTypeSpelling(ty);
+ std::string baseTypeName = getTypeSpelling(ty.getCanonicalType());
// Remove access qualifiers on images
// (as they are inseparable from type in clang implementation,
@@ -1380,38 +1649,13 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
}
argTypeNames.push_back(llvm::MDString::get(VMContext, typeName));
-
- // Turn "unsigned type" to "utype"
- pos = baseTypeName.find("unsigned");
- if (pos != std::string::npos)
- baseTypeName.erase(pos + 1, 8);
-
argBaseTypeNames.push_back(
llvm::MDString::get(VMContext, baseTypeName));
if (isPipe)
typeQuals = "pipe";
}
-
argTypeQuals.push_back(llvm::MDString::get(VMContext, typeQuals));
-
- // Get image and pipe access qualifier:
- if (ty->isImageType() || ty->isPipeType()) {
- const Decl *PDecl = parm;
- if (auto *TD = dyn_cast<TypedefType>(ty))
- PDecl = TD->getDecl();
- const OpenCLAccessAttr *A = PDecl->getAttr<OpenCLAccessAttr>();
- if (A && A->isWriteOnly())
- accessQuals.push_back(llvm::MDString::get(VMContext, "write_only"));
- else if (A && A->isReadWrite())
- accessQuals.push_back(llvm::MDString::get(VMContext, "read_write"));
- else
- accessQuals.push_back(llvm::MDString::get(VMContext, "read_only"));
- } else
- accessQuals.push_back(llvm::MDString::get(VMContext, "none"));
-
- // Get argument name.
- argNames.push_back(llvm::MDString::get(VMContext, parm->getName()));
}
Fn->setMetadata("kernel_arg_addr_space",
@@ -1484,6 +1728,9 @@ void CodeGenModule::SetLLVMFunctionAttributesForDefinition(const Decl *D,
if (CodeGenOpts.UnwindTables)
B.addAttribute(llvm::Attribute::UWTable);
+ if (CodeGenOpts.StackClashProtector)
+ B.addAttribute("probe-stack", "inline-asm");
+
if (!hasUnwindExceptions(LangOpts))
B.addAttribute(llvm::Attribute::NoUnwind);
@@ -1514,16 +1761,15 @@ void CodeGenModule::SetLLVMFunctionAttributesForDefinition(const Decl *D,
!CodeGenOpts.DisableO0ImplyOptNone && CodeGenOpts.OptimizationLevel == 0;
// We can't add optnone in the following cases, it won't pass the verifier.
ShouldAddOptNone &= !D->hasAttr<MinSizeAttr>();
- ShouldAddOptNone &= !F->hasFnAttribute(llvm::Attribute::AlwaysInline);
ShouldAddOptNone &= !D->hasAttr<AlwaysInlineAttr>();
- if (ShouldAddOptNone || D->hasAttr<OptimizeNoneAttr>()) {
+ // Add optnone, but do so only if the function isn't always_inline.
+ if ((ShouldAddOptNone || D->hasAttr<OptimizeNoneAttr>()) &&
+ !F->hasFnAttribute(llvm::Attribute::AlwaysInline)) {
B.addAttribute(llvm::Attribute::OptimizeNone);
// OptimizeNone implies noinline; we should not be inlining such functions.
B.addAttribute(llvm::Attribute::NoInline);
- assert(!F->hasFnAttribute(llvm::Attribute::AlwaysInline) &&
- "OptimizeNone and AlwaysInline on same function!");
// We still need to handle naked functions even though optnone subsumes
// much of their semantics.
@@ -1539,7 +1785,8 @@ void CodeGenModule::SetLLVMFunctionAttributesForDefinition(const Decl *D,
B.addAttribute(llvm::Attribute::NoInline);
} else if (D->hasAttr<NoDuplicateAttr>()) {
B.addAttribute(llvm::Attribute::NoDuplicate);
- } else if (D->hasAttr<NoInlineAttr>()) {
+ } else if (D->hasAttr<NoInlineAttr>() && !F->hasFnAttribute(llvm::Attribute::AlwaysInline)) {
+ // Add noinline if the function isn't always_inline.
B.addAttribute(llvm::Attribute::NoInline);
} else if (D->hasAttr<AlwaysInlineAttr>() &&
!F->hasFnAttribute(llvm::Attribute::NoInline)) {
@@ -1585,7 +1832,8 @@ void CodeGenModule::SetLLVMFunctionAttributesForDefinition(const Decl *D,
B.addAttribute(llvm::Attribute::OptimizeForSize);
B.addAttribute(llvm::Attribute::Cold);
}
-
+ if (D->hasAttr<HotAttr>())
+ B.addAttribute(llvm::Attribute::Hot);
if (D->hasAttr<MinSizeAttr>())
B.addAttribute(llvm::Attribute::MinSize);
}
@@ -1635,6 +1883,15 @@ void CodeGenModule::SetLLVMFunctionAttributesForDefinition(const Decl *D,
}
}
+void CodeGenModule::setLLVMFunctionFEnvAttributes(const FunctionDecl *D,
+ llvm::Function *F) {
+ if (D->hasAttr<StrictFPAttr>()) {
+ llvm::AttrBuilder FuncAttrs;
+ FuncAttrs.addAttribute("strictfp");
+ F->addAttributes(llvm::AttributeList::FunctionIndex, FuncAttrs);
+ }
+}
+
void CodeGenModule::SetCommonAttributes(GlobalDecl GD, llvm::GlobalValue *GV) {
const Decl *D = GD.getDecl();
if (dyn_cast_or_null<NamedDecl>(D))
@@ -1643,13 +1900,13 @@ void CodeGenModule::SetCommonAttributes(GlobalDecl GD, llvm::GlobalValue *GV) {
GV->setVisibility(llvm::GlobalValue::DefaultVisibility);
if (D && D->hasAttr<UsedAttr>())
- addUsedGlobal(GV);
+ addUsedOrCompilerUsedGlobal(GV);
if (CodeGenOpts.KeepStaticConsts && D && isa<VarDecl>(D)) {
const auto *VD = cast<VarDecl>(D);
if (VD->getType().isConstQualified() &&
VD->getStorageDuration() == SD_Static)
- addUsedGlobal(GV);
+ addUsedOrCompilerUsedGlobal(GV);
}
}
@@ -1659,6 +1916,7 @@ bool CodeGenModule::GetCPUAndFeaturesAttributes(GlobalDecl GD,
// we have a decl for the function and it has a target attribute then
// parse that and add it to the feature set.
StringRef TargetCPU = getTarget().getTargetOpts().CPU;
+ StringRef TuneCPU = getTarget().getTargetOpts().TuneCPU;
std::vector<std::string> Features;
const auto *FD = dyn_cast_or_null<FunctionDecl>(GD.getDecl());
FD = FD ? FD->getMostRecentDecl() : FD;
@@ -1667,7 +1925,7 @@ bool CodeGenModule::GetCPUAndFeaturesAttributes(GlobalDecl GD,
bool AddedAttr = false;
if (TD || SD) {
llvm::StringMap<bool> FeatureMap;
- getFunctionFeatureMap(FeatureMap, GD);
+ getContext().getFunctionFeatureMap(FeatureMap, GD);
// Produce the canonical string for this set of features.
for (const llvm::StringMap<bool>::value_type &Entry : FeatureMap)
@@ -1678,10 +1936,15 @@ bool CodeGenModule::GetCPUAndFeaturesAttributes(GlobalDecl GD,
// get and parse the target attribute so we can get the cpu for
// the function.
if (TD) {
- TargetAttr::ParsedTargetAttr ParsedAttr = TD->parse();
- if (ParsedAttr.Architecture != "" &&
- getTarget().isValidCPUName(ParsedAttr.Architecture))
+ ParsedTargetAttr ParsedAttr = TD->parse();
+ if (!ParsedAttr.Architecture.empty() &&
+ getTarget().isValidCPUName(ParsedAttr.Architecture)) {
TargetCPU = ParsedAttr.Architecture;
+ TuneCPU = ""; // Clear the tune CPU.
+ }
+ if (!ParsedAttr.Tune.empty() &&
+ getTarget().isValidCPUName(ParsedAttr.Tune))
+ TuneCPU = ParsedAttr.Tune;
}
} else {
// Otherwise just add the existing target cpu and target features to the
@@ -1689,10 +1952,14 @@ bool CodeGenModule::GetCPUAndFeaturesAttributes(GlobalDecl GD,
Features = getTarget().getTargetOpts().Features;
}
- if (TargetCPU != "") {
+ if (!TargetCPU.empty()) {
Attrs.addAttribute("target-cpu", TargetCPU);
AddedAttr = true;
}
+ if (!TuneCPU.empty()) {
+ Attrs.addAttribute("tune-cpu", TuneCPU);
+ AddedAttr = true;
+ }
if (!Features.empty()) {
llvm::sort(Features);
Attrs.addAttribute("target-features", llvm::join(Features, ","));
@@ -1709,6 +1976,8 @@ void CodeGenModule::setNonAliasAttributes(GlobalDecl GD,
if (D) {
if (auto *GV = dyn_cast<llvm::GlobalVariable>(GO)) {
+ if (D->hasAttr<RetainAttr>())
+ addUsedGlobal(GV);
if (auto *SA = D->getAttr<PragmaClangBSSSectionAttr>())
GV->addAttribute("bss-section", SA->getName());
if (auto *SA = D->getAttr<PragmaClangDataSectionAttr>())
@@ -1720,6 +1989,8 @@ void CodeGenModule::setNonAliasAttributes(GlobalDecl GD,
}
if (auto *F = dyn_cast<llvm::Function>(GO)) {
+ if (D->hasAttr<RetainAttr>())
+ addUsedGlobal(F);
if (auto *SA = D->getAttr<PragmaClangTextSectionAttr>())
if (!D->getAttr<SectionAttr>())
F->addFnAttr("implicit-section-name", SA->getName());
@@ -1729,8 +2000,11 @@ void CodeGenModule::setNonAliasAttributes(GlobalDecl GD,
// We know that GetCPUAndFeaturesAttributes will always have the
// newest set, since it has the newest possible FunctionDecl, so the
// new ones should replace the old.
- F->removeFnAttr("target-cpu");
- F->removeFnAttr("target-features");
+ llvm::AttrBuilder RemoveAttrs;
+ RemoveAttrs.addAttribute("target-cpu");
+ RemoveAttrs.addAttribute("target-features");
+ RemoveAttrs.addAttribute("tune-cpu");
+ F->removeAttributes(llvm::AttributeList::FunctionIndex, RemoveAttrs);
F->addAttributes(llvm::AttributeList::FunctionIndex, Attrs);
}
}
@@ -1748,7 +2022,7 @@ void CodeGenModule::SetInternalFunctionAttributes(GlobalDecl GD,
llvm::Function *F,
const CGFunctionInfo &FI) {
const Decl *D = GD.getDecl();
- SetLLVMFunctionAttributes(GD, FI, F);
+ SetLLVMFunctionAttributes(GD, FI, F, /*IsThunk=*/false);
SetLLVMFunctionAttributesForDefinition(D, F);
F->setLinkage(llvm::Function::InternalLinkage);
@@ -1802,7 +2076,8 @@ void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F,
const auto *FD = cast<FunctionDecl>(GD.getDecl());
if (!IsIncompleteFunction)
- SetLLVMFunctionAttributes(GD, getTypes().arrangeGlobalDeclaration(GD), F);
+ SetLLVMFunctionAttributes(GD, getTypes().arrangeGlobalDeclaration(GD), F,
+ IsThunk);
// Add the Returned attribute for "this", except for iOS 5 and earlier
// where substantial code, including the libstdc++ dylib, was compiled with
@@ -1831,20 +2106,23 @@ void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F,
else if (const auto *SA = FD->getAttr<SectionAttr>())
F->setSection(SA->getName());
+ // If we plan on emitting this inline builtin, we can't treat it as a builtin.
+ if (FD->isInlineBuiltinDeclaration()) {
+ const FunctionDecl *FDBody;
+ bool HasBody = FD->hasBody(FDBody);
+ (void)HasBody;
+ assert(HasBody && "Inline builtin declarations should always have an "
+ "available body!");
+ if (shouldEmitFunction(FDBody))
+ F->addAttribute(llvm::AttributeList::FunctionIndex,
+ llvm::Attribute::NoBuiltin);
+ }
+
if (FD->isReplaceableGlobalAllocationFunction()) {
// A replaceable global allocation function does not act like a builtin by
// default, only if it is invoked by a new-expression or delete-expression.
F->addAttribute(llvm::AttributeList::FunctionIndex,
llvm::Attribute::NoBuiltin);
-
- // A sane operator new returns a non-aliasing pointer.
- // FIXME: Also add NonNull attribute to the return value
- // for the non-nothrow forms?
- auto Kind = FD->getDeclName().getCXXOverloadedOperator();
- if (getCodeGenOpts().AssumeSaneOperatorNew &&
- (Kind == OO_New || Kind == OO_Array_New))
- F->addAttribute(llvm::AttributeList::ReturnIndex,
- llvm::Attribute::NoAlias);
}
if (isa<CXXConstructorDecl>(FD) || isa<CXXDestructorDecl>(FD))
@@ -1883,7 +2161,7 @@ void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F,
}
void CodeGenModule::addUsedGlobal(llvm::GlobalValue *GV) {
- assert(!GV->isDeclaration() &&
+ assert((isa<llvm::Function>(GV) || !GV->isDeclaration()) &&
"Only globals with definition can force usage.");
LLVMUsed.emplace_back(GV);
}
@@ -1894,6 +2172,15 @@ void CodeGenModule::addCompilerUsedGlobal(llvm::GlobalValue *GV) {
LLVMCompilerUsed.emplace_back(GV);
}
+void CodeGenModule::addUsedOrCompilerUsedGlobal(llvm::GlobalValue *GV) {
+ assert((isa<llvm::Function>(GV) || !GV->isDeclaration()) &&
+ "Only globals with definition can force usage.");
+ if (getTriple().isOSBinFormatELF())
+ LLVMCompilerUsed.emplace_back(GV);
+ else
+ LLVMUsed.emplace_back(GV);
+}
+
static void emitUsed(CodeGenModule &CGM, StringRef Name,
std::vector<llvm::WeakTrackingVH> &List) {
// Don't create llvm.used if there is no need.
@@ -1933,6 +2220,8 @@ void CodeGenModule::AppendLinkerOptions(StringRef Opts) {
void CodeGenModule::AddDetectMismatch(StringRef Name, StringRef Value) {
llvm::SmallString<32> Opt;
getTargetCodeGenInfo().getDetectMismatchOption(Name, Value, Opt);
+ if (Opt.empty())
+ return;
auto *MDOpts = llvm::MDString::get(getLLVMContext(), Opt);
LinkerOptionsMetadata.push_back(llvm::MDNode::get(getLLVMContext(), MDOpts));
}
@@ -2087,6 +2376,13 @@ void CodeGenModule::EmitDeferred() {
assert(DeferredVTables.empty());
}
+ // Emit CUDA/HIP static device variables referenced by host code only.
+ // Note we should not clear CUDADeviceVarODRUsedByHost since it is still
+ // needed for further handling.
+ if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice)
+ for (const auto *V : getContext().CUDADeviceVarODRUsedByHost)
+ DeferredDeclsToEmit.push_back(V);
+
// Stop if we're out of both deferred vtables and deferred declarations.
if (DeferredDeclsToEmit.empty())
return;
@@ -2202,20 +2498,61 @@ llvm::Constant *CodeGenModule::EmitAnnotationLineNo(SourceLocation L) {
return llvm::ConstantInt::get(Int32Ty, LineNo);
}
+llvm::Constant *CodeGenModule::EmitAnnotationArgs(const AnnotateAttr *Attr) {
+ ArrayRef<Expr *> Exprs = {Attr->args_begin(), Attr->args_size()};
+ if (Exprs.empty())
+ return llvm::ConstantPointerNull::get(Int8PtrTy);
+
+ llvm::FoldingSetNodeID ID;
+ for (Expr *E : Exprs) {
+ ID.Add(cast<clang::ConstantExpr>(E)->getAPValueResult());
+ }
+ llvm::Constant *&Lookup = AnnotationArgs[ID.ComputeHash()];
+ if (Lookup)
+ return Lookup;
+
+ llvm::SmallVector<llvm::Constant *, 4> LLVMArgs;
+ LLVMArgs.reserve(Exprs.size());
+ ConstantEmitter ConstEmiter(*this);
+ llvm::transform(Exprs, std::back_inserter(LLVMArgs), [&](const Expr *E) {
+ const auto *CE = cast<clang::ConstantExpr>(E);
+ return ConstEmiter.emitAbstract(CE->getBeginLoc(), CE->getAPValueResult(),
+ CE->getType());
+ });
+ auto *Struct = llvm::ConstantStruct::getAnon(LLVMArgs);
+ auto *GV = new llvm::GlobalVariable(getModule(), Struct->getType(), true,
+ llvm::GlobalValue::PrivateLinkage, Struct,
+ ".args");
+ GV->setSection(AnnotationSection);
+ GV->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
+ auto *Bitcasted = llvm::ConstantExpr::getBitCast(GV, Int8PtrTy);
+
+ Lookup = Bitcasted;
+ return Bitcasted;
+}
+
llvm::Constant *CodeGenModule::EmitAnnotateAttr(llvm::GlobalValue *GV,
const AnnotateAttr *AA,
SourceLocation L) {
// Get the globals for file name, annotation, and the line number.
llvm::Constant *AnnoGV = EmitAnnotationString(AA->getAnnotation()),
*UnitGV = EmitAnnotationUnit(L),
- *LineNoCst = EmitAnnotationLineNo(L);
+ *LineNoCst = EmitAnnotationLineNo(L),
+ *Args = EmitAnnotationArgs(AA);
+
+ llvm::Constant *ASZeroGV = GV;
+ if (GV->getAddressSpace() != 0) {
+ ASZeroGV = llvm::ConstantExpr::getAddrSpaceCast(
+ GV, GV->getValueType()->getPointerTo(0));
+ }
// Create the ConstantStruct for the global annotation.
- llvm::Constant *Fields[4] = {
- llvm::ConstantExpr::getBitCast(GV, Int8PtrTy),
- llvm::ConstantExpr::getBitCast(AnnoGV, Int8PtrTy),
- llvm::ConstantExpr::getBitCast(UnitGV, Int8PtrTy),
- LineNoCst
+ llvm::Constant *Fields[] = {
+ llvm::ConstantExpr::getBitCast(ASZeroGV, Int8PtrTy),
+ llvm::ConstantExpr::getBitCast(AnnoGV, Int8PtrTy),
+ llvm::ConstantExpr::getBitCast(UnitGV, Int8PtrTy),
+ LineNoCst,
+ Args,
};
return llvm::ConstantStruct::getAnon(Fields);
}
@@ -2228,29 +2565,28 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D,
Annotations.push_back(EmitAnnotateAttr(GV, I, D->getLocation()));
}
-bool CodeGenModule::isInSanitizerBlacklist(SanitizerMask Kind,
- llvm::Function *Fn,
- SourceLocation Loc) const {
- const auto &SanitizerBL = getContext().getSanitizerBlacklist();
- // Blacklist by function name.
- if (SanitizerBL.isBlacklistedFunction(Kind, Fn->getName()))
+bool CodeGenModule::isInNoSanitizeList(SanitizerMask Kind, llvm::Function *Fn,
+ SourceLocation Loc) const {
+ const auto &NoSanitizeL = getContext().getNoSanitizeList();
+ // NoSanitize by function name.
+ if (NoSanitizeL.containsFunction(Kind, Fn->getName()))
return true;
- // Blacklist by location.
+ // NoSanitize by location.
if (Loc.isValid())
- return SanitizerBL.isBlacklistedLocation(Kind, Loc);
+ return NoSanitizeL.containsLocation(Kind, Loc);
// If location is unknown, this may be a compiler-generated function. Assume
// it's located in the main file.
auto &SM = Context.getSourceManager();
if (const auto *MainFile = SM.getFileEntryForID(SM.getMainFileID())) {
- return SanitizerBL.isBlacklistedFile(Kind, MainFile->getName());
+ return NoSanitizeL.containsFile(Kind, MainFile->getName());
}
return false;
}
-bool CodeGenModule::isInSanitizerBlacklist(llvm::GlobalVariable *GV,
- SourceLocation Loc, QualType Ty,
- StringRef Category) const {
- // For now globals can be blacklisted only in ASan and KASan.
+bool CodeGenModule::isInNoSanitizeList(llvm::GlobalVariable *GV,
+ SourceLocation Loc, QualType Ty,
+ StringRef Category) const {
+ // For now globals can be ignored only in ASan and KASan.
const SanitizerMask EnabledAsanMask =
LangOpts.Sanitize.Mask &
(SanitizerKind::Address | SanitizerKind::KernelAddress |
@@ -2258,22 +2594,22 @@ bool CodeGenModule::isInSanitizerBlacklist(llvm::GlobalVariable *GV,
SanitizerKind::MemTag);
if (!EnabledAsanMask)
return false;
- const auto &SanitizerBL = getContext().getSanitizerBlacklist();
- if (SanitizerBL.isBlacklistedGlobal(EnabledAsanMask, GV->getName(), Category))
+ const auto &NoSanitizeL = getContext().getNoSanitizeList();
+ if (NoSanitizeL.containsGlobal(EnabledAsanMask, GV->getName(), Category))
return true;
- if (SanitizerBL.isBlacklistedLocation(EnabledAsanMask, Loc, Category))
+ if (NoSanitizeL.containsLocation(EnabledAsanMask, Loc, Category))
return true;
// Check global type.
if (!Ty.isNull()) {
// Drill down the array types: if global variable of a fixed type is
- // blacklisted, we also don't instrument arrays of them.
+ // not sanitized, we also don't instrument arrays of them.
while (auto AT = dyn_cast<ArrayType>(Ty.getTypePtr()))
Ty = AT->getElementType();
Ty = Ty.getCanonicalType().getUnqualifiedType();
- // We allow to blacklist only record types (classes, structs etc.)
+ // Only record types (classes, structs etc.) are ignored.
if (Ty->isRecordType()) {
std::string TypeStr = Ty.getAsString(getContext().getPrintingPolicy());
- if (SanitizerBL.isBlacklistedType(EnabledAsanMask, TypeStr, Category))
+ if (NoSanitizeL.containsType(EnabledAsanMask, TypeStr, Category))
return true;
}
}
@@ -2306,6 +2642,34 @@ bool CodeGenModule::imbueXRayAttrs(llvm::Function *Fn, SourceLocation Loc,
return true;
}
+bool CodeGenModule::isProfileInstrExcluded(llvm::Function *Fn,
+ SourceLocation Loc) const {
+ const auto &ProfileList = getContext().getProfileList();
+ // If the profile list is empty, then instrument everything.
+ if (ProfileList.isEmpty())
+ return false;
+ CodeGenOptions::ProfileInstrKind Kind = getCodeGenOpts().getProfileInstr();
+ // First, check the function name.
+ Optional<bool> V = ProfileList.isFunctionExcluded(Fn->getName(), Kind);
+ if (V.hasValue())
+ return *V;
+ // Next, check the source location.
+ if (Loc.isValid()) {
+ Optional<bool> V = ProfileList.isLocationExcluded(Loc, Kind);
+ if (V.hasValue())
+ return *V;
+ }
+ // If location is unknown, this may be a compiler-generated function. Assume
+ // it's located in the main file.
+ auto &SM = Context.getSourceManager();
+ if (const auto *MainFile = SM.getFileEntryForID(SM.getMainFileID())) {
+ Optional<bool> V = ProfileList.isFileExcluded(MainFile->getName(), Kind);
+ if (V.hasValue())
+ return *V;
+ }
+ return ProfileList.getDefault();
+}
+
bool CodeGenModule::MustBeEmitted(const ValueDecl *Global) {
// Never defer when EmitAllDecls is specified.
if (LangOpts.EmitAllDecls)
@@ -2322,19 +2686,24 @@ bool CodeGenModule::MustBeEmitted(const ValueDecl *Global) {
}
bool CodeGenModule::MayBeEmittedEagerly(const ValueDecl *Global) {
+ // In OpenMP 5.0 variables and function may be marked as
+ // device_type(host/nohost) and we should not emit them eagerly unless we sure
+ // that they must be emitted on the host/device. To be sure we need to have
+ // seen a declare target with an explicit mentioning of the function, we know
+ // we have if the level of the declare target attribute is -1. Note that we
+ // check somewhere else if we should emit this at all.
+ if (LangOpts.OpenMP >= 50 && !LangOpts.OpenMPSimd) {
+ llvm::Optional<OMPDeclareTargetDeclAttr *> ActiveAttr =
+ OMPDeclareTargetDeclAttr::getActiveAttr(Global);
+ if (!ActiveAttr || (*ActiveAttr)->getLevel() != (unsigned)-1)
+ return false;
+ }
+
if (const auto *FD = dyn_cast<FunctionDecl>(Global)) {
if (FD->getTemplateSpecializationKind() == TSK_ImplicitInstantiation)
// Implicit template instantiations may change linkage if they are later
// explicitly instantiated, so they should not be emitted eagerly.
return false;
- // In OpenMP 5.0 function may be marked as device_type(nohost) and we should
- // not emit them eagerly unless we sure that the function must be emitted on
- // the host.
- if (LangOpts.OpenMP >= 50 && !LangOpts.OpenMPSimd &&
- !LangOpts.OpenMPIsDevice &&
- !OMPDeclareTargetDeclAttr::getDeviceType(FD) &&
- !FD->isUsed(/*CheckUsedAttr=*/false) && !FD->isReferenced())
- return false;
}
if (const auto *VD = dyn_cast<VarDecl>(Global))
if (Context.getInlineVariableDefinitionKind(VD) ==
@@ -2353,13 +2722,8 @@ bool CodeGenModule::MayBeEmittedEagerly(const ValueDecl *Global) {
return true;
}
-ConstantAddress CodeGenModule::GetAddrOfUuidDescriptor(
- const CXXUuidofExpr* E) {
- // Sema has verified that IIDSource has a __declspec(uuid()), and that its
- // well-formed.
- StringRef Uuid = E->getUuidStr();
- std::string Name = "_GUID_" + Uuid.lower();
- std::replace(Name.begin(), Name.end(), '-', '_');
+ConstantAddress CodeGenModule::GetAddrOfMSGuidDecl(const MSGuidDecl *GD) {
+ StringRef Name = getMangledName(GD);
// The UUID descriptor should be pointer aligned.
CharUnits Alignment = CharUnits::fromQuantity(PointerAlignInBytes);
@@ -2368,8 +2732,30 @@ ConstantAddress CodeGenModule::GetAddrOfUuidDescriptor(
if (llvm::GlobalVariable *GV = getModule().getNamedGlobal(Name))
return ConstantAddress(GV, Alignment);
- llvm::Constant *Init = EmitUuidofInitializer(Uuid);
- assert(Init && "failed to initialize as constant");
+ ConstantEmitter Emitter(*this);
+ llvm::Constant *Init;
+
+ APValue &V = GD->getAsAPValue();
+ if (!V.isAbsent()) {
+ // If possible, emit the APValue version of the initializer. In particular,
+ // this gets the type of the constant right.
+ Init = Emitter.emitForInitializer(
+ GD->getAsAPValue(), GD->getType().getAddressSpace(), GD->getType());
+ } else {
+ // As a fallback, directly construct the constant.
+ // FIXME: This may get padding wrong under esoteric struct layout rules.
+ // MSVC appears to create a complete type 'struct __s_GUID' that it
+ // presumably uses to represent these constants.
+ MSGuidDecl::Parts Parts = GD->getParts();
+ llvm::Constant *Fields[4] = {
+ llvm::ConstantInt::get(Int32Ty, Parts.Part1),
+ llvm::ConstantInt::get(Int16Ty, Parts.Part2),
+ llvm::ConstantInt::get(Int16Ty, Parts.Part3),
+ llvm::ConstantDataArray::getRaw(
+ StringRef(reinterpret_cast<char *>(Parts.Part4And5), 8), 8,
+ Int8Ty)};
+ Init = llvm::ConstantStruct::getAnon(Fields);
+ }
auto *GV = new llvm::GlobalVariable(
getModule(), Init->getType(),
@@ -2377,6 +2763,42 @@ ConstantAddress CodeGenModule::GetAddrOfUuidDescriptor(
if (supportsCOMDAT())
GV->setComdat(TheModule.getOrInsertComdat(GV->getName()));
setDSOLocal(GV);
+
+ llvm::Constant *Addr = GV;
+ if (!V.isAbsent()) {
+ Emitter.finalize(GV);
+ } else {
+ llvm::Type *Ty = getTypes().ConvertTypeForMem(GD->getType());
+ Addr = llvm::ConstantExpr::getBitCast(
+ GV, Ty->getPointerTo(GV->getAddressSpace()));
+ }
+ return ConstantAddress(Addr, Alignment);
+}
+
+ConstantAddress CodeGenModule::GetAddrOfTemplateParamObject(
+ const TemplateParamObjectDecl *TPO) {
+ StringRef Name = getMangledName(TPO);
+ CharUnits Alignment = getNaturalTypeAlignment(TPO->getType());
+
+ if (llvm::GlobalVariable *GV = getModule().getNamedGlobal(Name))
+ return ConstantAddress(GV, Alignment);
+
+ ConstantEmitter Emitter(*this);
+ llvm::Constant *Init = Emitter.emitForInitializer(
+ TPO->getValue(), TPO->getType().getAddressSpace(), TPO->getType());
+
+ if (!Init) {
+ ErrorUnsupported(TPO, "template parameter object");
+ return ConstantAddress::invalid();
+ }
+
+ auto *GV = new llvm::GlobalVariable(
+ getModule(), Init->getType(),
+ /*isConstant=*/true, llvm::GlobalValue::LinkOnceODRLinkage, Init, Name);
+ if (supportsCOMDAT())
+ GV->setComdat(TheModule.getOrInsertComdat(GV->getName()));
+ Emitter.finalize(GV);
+
return ConstantAddress(GV, Alignment);
}
@@ -2401,9 +2823,7 @@ ConstantAddress CodeGenModule::GetWeakRefReference(const ValueDecl *VD) {
GlobalDecl(cast<FunctionDecl>(VD)),
/*ForVTable=*/false);
else
- Aliasee = GetOrCreateLLVMGlobal(AA->getAliasee(),
- llvm::PointerType::getUnqual(DeclTy),
- nullptr);
+ Aliasee = GetOrCreateLLVMGlobal(AA->getAliasee(), DeclTy, 0, nullptr);
auto *F = cast<llvm::GlobalValue>(Aliasee);
F->setLinkage(llvm::Function::ExternalWeakLinkage);
@@ -2439,7 +2859,8 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
!Global->hasAttr<CUDAGlobalAttr>() &&
!Global->hasAttr<CUDAConstantAttr>() &&
!Global->hasAttr<CUDASharedAttr>() &&
- !(LangOpts.HIP && Global->hasAttr<HIPPinnedShadowAttr>()))
+ !Global->getType()->isCUDADeviceBuiltinSurfaceType() &&
+ !Global->getType()->isCUDADeviceBuiltinTextureType())
return;
} else {
// We need to emit host-side 'shadows' for all global
@@ -2532,11 +2953,6 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
return;
}
- // Check if this must be emitted as declare variant.
- if (LangOpts.OpenMP && isa<FunctionDecl>(Global) && OpenMPRuntime &&
- OpenMPRuntime->emitDeclareVariant(GD, /*IsForDefinition=*/false))
- return;
-
// If we're deferring emission of a C++ variable with an
// initializer, remember the order in which it appeared in the file.
if (getLangOpts().CPlusPlus && isa<VarDecl>(Global) &&
@@ -2697,7 +3113,7 @@ bool CodeGenModule::shouldEmitFunction(GlobalDecl GD) {
if (CodeGenOpts.OptimizationLevel == 0 && !F->hasAttr<AlwaysInlineAttr>())
return false;
- if (F->hasAttr<DLLImportAttr>()) {
+ if (F->hasAttr<DLLImportAttr>() && !F->hasAttr<AlwaysInlineAttr>()) {
// Check whether it would be safe to inline this dllimport function.
DLLImportFunctionVisitor Visitor;
Visitor.TraverseFunctionDecl(const_cast<FunctionDecl*>(F));
@@ -2719,8 +3135,8 @@ bool CodeGenModule::shouldEmitFunction(GlobalDecl GD) {
// PR9614. Avoid cases where the source code is lying to us. An available
// externally function should have an equivalent function somewhere else,
- // but a function that calls itself is clearly not equivalent to the real
- // implementation.
+ // but a function that calls itself through asm label/`__builtin_` trickery is
+ // clearly not equivalent to the real implementation.
// This happens in glibc's btowc and in some configure checks.
return !isTriviallyRecursive(F);
}
@@ -2742,50 +3158,6 @@ void CodeGenModule::EmitMultiVersionFunctionDefinition(GlobalDecl GD,
EmitGlobalFunctionDefinition(GD, GV);
}
-void CodeGenModule::emitOpenMPDeviceFunctionRedefinition(
- GlobalDecl OldGD, GlobalDecl NewGD, llvm::GlobalValue *GV) {
- assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
- OpenMPRuntime && "Expected OpenMP device mode.");
- const auto *D = cast<FunctionDecl>(OldGD.getDecl());
-
- // Compute the function info and LLVM type.
- const CGFunctionInfo &FI = getTypes().arrangeGlobalDeclaration(OldGD);
- llvm::FunctionType *Ty = getTypes().GetFunctionType(FI);
-
- // Get or create the prototype for the function.
- if (!GV || (GV->getType()->getElementType() != Ty)) {
- GV = cast<llvm::GlobalValue>(GetOrCreateLLVMFunction(
- getMangledName(OldGD), Ty, GlobalDecl(), /*ForVTable=*/false,
- /*DontDefer=*/true, /*IsThunk=*/false, llvm::AttributeList(),
- ForDefinition));
- SetFunctionAttributes(OldGD, cast<llvm::Function>(GV),
- /*IsIncompleteFunction=*/false,
- /*IsThunk=*/false);
- }
- // We need to set linkage and visibility on the function before
- // generating code for it because various parts of IR generation
- // want to propagate this information down (e.g. to local static
- // declarations).
- auto *Fn = cast<llvm::Function>(GV);
- setFunctionLinkage(OldGD, Fn);
-
- // FIXME: this is redundant with part of
- // setFunctionDefinitionAttributes
- setGVProperties(Fn, OldGD);
-
- MaybeHandleStaticInExternC(D, Fn);
-
- maybeSetTrivialComdat(*D, *Fn);
-
- CodeGenFunction(*this).GenerateCode(NewGD, Fn, FI);
-
- setNonAliasAttributes(OldGD, Fn);
- SetLLVMFunctionAttributesForDefinition(D, Fn);
-
- if (D->hasAttr<AnnotateAttr>())
- AddGlobalAnnotations(D, Fn);
-}
-
void CodeGenModule::EmitGlobalDefinition(GlobalDecl GD, llvm::GlobalValue *GV) {
const auto *D = cast<ValueDecl>(GD.getDecl());
@@ -2851,7 +3223,9 @@ TargetMVPriority(const TargetInfo &TI,
}
void CodeGenModule::emitMultiVersionFunctions() {
- for (GlobalDecl GD : MultiVersionFuncs) {
+ std::vector<GlobalDecl> MVFuncsToEmit;
+ MultiVersionFuncs.swap(MVFuncsToEmit);
+ for (GlobalDecl GD : MVFuncsToEmit) {
SmallVector<CodeGenFunction::MultiVersionResolverOption, 10> Options;
const FunctionDecl *FD = cast<FunctionDecl>(GD.getDecl());
getContext().forEachMultiversionedFunctionVersion(
@@ -2905,6 +3279,17 @@ void CodeGenModule::emitMultiVersionFunctions() {
CodeGenFunction CGF(*this);
CGF.EmitMultiVersionResolver(ResolverFunc, Options);
}
+
+ // Ensure that any additions to the deferred decls list caused by emitting a
+ // variant are emitted. This can happen when the variant itself is inline and
+ // calls a function without linkage.
+ if (!MVFuncsToEmit.empty())
+ EmitDeferred();
+
+ // Ensure that any additions to the multiversion funcs list from either the
+ // deferred decls or the multiversion functions themselves are emitted.
+ if (!MultiVersionFuncs.empty())
+ emitMultiVersionFunctions();
}
void CodeGenModule::emitCPUDispatchDefinition(GlobalDecl GD) {
@@ -2979,7 +3364,7 @@ void CodeGenModule::emitCPUDispatchDefinition(GlobalDecl GD) {
++Index;
}
- llvm::sort(
+ llvm::stable_sort(
Options, [](const CodeGenFunction::MultiVersionResolverOption &LHS,
const CodeGenFunction::MultiVersionResolverOption &RHS) {
return CodeGenFunction::GetX86CpuSupportsMask(LHS.Conditions.Features) >
@@ -3100,14 +3485,9 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction(
EmitGlobal(GDDef);
}
}
- // Check if this must be emitted as declare variant and emit reference to
- // the the declare variant function.
- if (LangOpts.OpenMP && OpenMPRuntime)
- (void)OpenMPRuntime->emitDeclareVariant(GD, /*IsForDefinition=*/true);
if (FD->isMultiVersion()) {
- const auto *TA = FD->getAttr<TargetAttr>();
- if (TA && TA->isDefaultVersion())
+ if (FD->hasAttr<TargetAttr>())
UpdateMultiVersionNames(GD, FD);
if (!IsForDefinition)
return GetOrCreateMultiVersionResolver(GD, Ty, FD);
@@ -3147,7 +3527,7 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction(
}
if ((isa<llvm::Function>(Entry) || isa<llvm::GlobalAlias>(Entry)) &&
- (Entry->getType()->getElementType() == Ty)) {
+ (Entry->getValueType() == Ty)) {
return Entry;
}
@@ -3196,7 +3576,7 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction(
}
llvm::Constant *BC = llvm::ConstantExpr::getBitCast(
- F, Entry->getType()->getElementType()->getPointerTo());
+ F, Entry->getValueType()->getPointerTo());
addGlobalValReplacement(Entry, BC);
}
@@ -3255,7 +3635,7 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction(
// Make sure the result is of the requested type.
if (!IsIncompleteFunction) {
- assert(F->getType()->getElementType() == Ty);
+ assert(F->getFunctionType() == Ty);
return F;
}
@@ -3271,6 +3651,8 @@ llvm::Constant *CodeGenModule::GetAddrOfFunction(GlobalDecl GD,
bool ForVTable,
bool DontDefer,
ForDefinition_t IsForDefinition) {
+ assert(!cast<FunctionDecl>(GD.getDecl())->isConsteval() &&
+ "consteval function should never be emitted");
// If there was no specific requested type, just convert it now.
if (!Ty) {
const auto *FD = cast<FunctionDecl>(GD.getDecl());
@@ -3288,9 +3670,19 @@ llvm::Constant *CodeGenModule::GetAddrOfFunction(GlobalDecl GD,
}
StringRef MangledName = getMangledName(GD);
- return GetOrCreateLLVMFunction(MangledName, Ty, GD, ForVTable, DontDefer,
- /*IsThunk=*/false, llvm::AttributeList(),
- IsForDefinition);
+ auto *F = GetOrCreateLLVMFunction(MangledName, Ty, GD, ForVTable, DontDefer,
+ /*IsThunk=*/false, llvm::AttributeList(),
+ IsForDefinition);
+ // Returns kernel handle for HIP kernel stub function.
+ if (LangOpts.CUDA && !LangOpts.CUDAIsDevice &&
+ cast<FunctionDecl>(GD.getDecl())->hasAttr<CUDAGlobalAttr>()) {
+ auto *Handle = getCUDARuntime().getKernelHandle(
+ cast<llvm::Function>(F->stripPointerCasts()), GD);
+ if (IsForDefinition)
+ return F;
+ return llvm::ConstantExpr::getBitCast(Handle, Ty->getPointerTo());
+ }
+ return F;
}
static const FunctionDecl *
@@ -3299,8 +3691,8 @@ GetRuntimeFunctionDecl(ASTContext &C, StringRef Name) {
DeclContext *DC = TranslationUnitDecl::castToDeclContext(TUDecl);
IdentifierInfo &CII = C.Idents.get(Name);
- for (const auto &Result : DC->lookup(&CII))
- if (const auto FD = dyn_cast<FunctionDecl>(Result))
+ for (const auto *Result : DC->lookup(&CII))
+ if (const auto *FD = dyn_cast<FunctionDecl>(Result))
return FD;
if (!C.getLangOpts().CPlusPlus)
@@ -3314,15 +3706,15 @@ GetRuntimeFunctionDecl(ASTContext &C, StringRef Name) {
for (const auto &N : {"__cxxabiv1", "std"}) {
IdentifierInfo &NS = C.Idents.get(N);
- for (const auto &Result : DC->lookup(&NS)) {
- NamespaceDecl *ND = dyn_cast<NamespaceDecl>(Result);
- if (auto LSD = dyn_cast<LinkageSpecDecl>(Result))
- for (const auto &Result : LSD->lookup(&NS))
+ for (const auto *Result : DC->lookup(&NS)) {
+ const NamespaceDecl *ND = dyn_cast<NamespaceDecl>(Result);
+ if (auto *LSD = dyn_cast<LinkageSpecDecl>(Result))
+ for (const auto *Result : LSD->lookup(&NS))
if ((ND = dyn_cast<NamespaceDecl>(Result)))
break;
if (ND)
- for (const auto &Result : ND->lookup(&CXXII))
+ for (const auto *Result : ND->lookup(&CXXII))
if (const auto *FD = dyn_cast<FunctionDecl>(Result))
return FD;
}
@@ -3393,9 +3785,9 @@ bool CodeGenModule::isTypeConstant(QualType Ty, bool ExcludeCtor) {
}
/// GetOrCreateLLVMGlobal - If the specified mangled name is not in the module,
-/// create and return an llvm GlobalVariable with the specified type. If there
-/// is something in the module with the specified name, return it potentially
-/// bitcasted to the right type.
+/// create and return an llvm GlobalVariable with the specified type and address
+/// space. If there is something in the module with the specified name, return
+/// it potentially bitcasted to the right type.
///
/// If D is non-null, it specifies a decl that correspond to this. This is used
/// to set the attributes on the global when it is first created.
@@ -3404,9 +3796,8 @@ bool CodeGenModule::isTypeConstant(QualType Ty, bool ExcludeCtor) {
/// type Ty will be returned, not conversion of a variable with the same
/// mangled name but some other type.
llvm::Constant *
-CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName,
- llvm::PointerType *Ty,
- const VarDecl *D,
+CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName, llvm::Type *Ty,
+ unsigned AddrSpace, const VarDecl *D,
ForDefinition_t IsForDefinition) {
// Lookup the entry, lazily creating it if necessary.
llvm::GlobalValue *Entry = GetGlobalValue(MangledName);
@@ -3423,7 +3814,7 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName,
if (LangOpts.OpenMP && !LangOpts.OpenMPSimd && D)
getOpenMPRuntime().registerTargetGlobalVariable(D, Entry);
- if (Entry->getType() == Ty)
+ if (Entry->getValueType() == Ty && Entry->getAddressSpace() == AddrSpace)
return Entry;
// If there are two attempts to define the same mangled name, issue an
@@ -3447,22 +3838,24 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName,
}
// Make sure the result is of the correct type.
- if (Entry->getType()->getAddressSpace() != Ty->getAddressSpace())
- return llvm::ConstantExpr::getAddrSpaceCast(Entry, Ty);
+ if (Entry->getType()->getAddressSpace() != AddrSpace) {
+ return llvm::ConstantExpr::getAddrSpaceCast(Entry,
+ Ty->getPointerTo(AddrSpace));
+ }
// (If global is requested for a definition, we always need to create a new
// global, not just return a bitcast.)
if (!IsForDefinition)
- return llvm::ConstantExpr::getBitCast(Entry, Ty);
+ return llvm::ConstantExpr::getBitCast(Entry, Ty->getPointerTo(AddrSpace));
}
- auto AddrSpace = GetGlobalVarAddressSpace(D);
- auto TargetAddrSpace = getContext().getTargetAddressSpace(AddrSpace);
+ auto DAddrSpace = GetGlobalVarAddressSpace(D);
+ auto TargetAddrSpace = getContext().getTargetAddressSpace(DAddrSpace);
auto *GV = new llvm::GlobalVariable(
- getModule(), Ty->getElementType(), false,
- llvm::GlobalValue::ExternalLinkage, nullptr, MangledName, nullptr,
- llvm::GlobalVariable::NotThreadLocal, TargetAddrSpace);
+ getModule(), Ty, false, llvm::GlobalValue::ExternalLinkage, nullptr,
+ MangledName, nullptr, llvm::GlobalVariable::NotThreadLocal,
+ TargetAddrSpace);
// If we already created a global with the same mangled name (but different
// type) before, take its name and remove it from its parent.
@@ -3546,7 +3939,7 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName,
llvm::Constant *Init = emitter.tryEmitForInitializer(*InitDecl);
if (Init) {
auto *InitType = Init->getType();
- if (GV->getType()->getElementType() != InitType) {
+ if (GV->getValueType() != InitType) {
// The type of the initializer does not match the definition.
// This happens when an initializer has a different type from
// the type of the global (because of padding at the end of a
@@ -3573,42 +3966,51 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName,
}
}
+ if (GV->isDeclaration()) {
+ getTargetCodeGenInfo().setTargetAttributes(D, GV, *this);
+ // External HIP managed variables needed to be recorded for transformation
+ // in both device and host compilations.
+ if (getLangOpts().CUDA && D && D->hasAttr<HIPManagedAttr>() &&
+ D->hasExternalStorage())
+ getCUDARuntime().handleVarRegistration(D, *GV);
+ }
+
LangAS ExpectedAS =
D ? D->getType().getAddressSpace()
: (LangOpts.OpenCL ? LangAS::opencl_global : LangAS::Default);
- assert(getContext().getTargetAddressSpace(ExpectedAS) ==
- Ty->getPointerAddressSpace());
- if (AddrSpace != ExpectedAS)
- return getTargetCodeGenInfo().performAddrSpaceCast(*this, GV, AddrSpace,
- ExpectedAS, Ty);
-
- if (GV->isDeclaration())
- getTargetCodeGenInfo().setTargetAttributes(D, GV, *this);
+ assert(getContext().getTargetAddressSpace(ExpectedAS) == AddrSpace);
+ if (DAddrSpace != ExpectedAS) {
+ return getTargetCodeGenInfo().performAddrSpaceCast(
+ *this, GV, DAddrSpace, ExpectedAS, Ty->getPointerTo(AddrSpace));
+ }
return GV;
}
llvm::Constant *
-CodeGenModule::GetAddrOfGlobal(GlobalDecl GD,
- ForDefinition_t IsForDefinition) {
+CodeGenModule::GetAddrOfGlobal(GlobalDecl GD, ForDefinition_t IsForDefinition) {
const Decl *D = GD.getDecl();
+
if (isa<CXXConstructorDecl>(D) || isa<CXXDestructorDecl>(D))
return getAddrOfCXXStructor(GD, /*FnInfo=*/nullptr, /*FnType=*/nullptr,
/*DontDefer=*/false, IsForDefinition);
- else if (isa<CXXMethodDecl>(D)) {
- auto FInfo = &getTypes().arrangeCXXMethodDeclaration(
- cast<CXXMethodDecl>(D));
+
+ if (isa<CXXMethodDecl>(D)) {
+ auto FInfo =
+ &getTypes().arrangeCXXMethodDeclaration(cast<CXXMethodDecl>(D));
auto Ty = getTypes().GetFunctionType(*FInfo);
return GetAddrOfFunction(GD, Ty, /*ForVTable=*/false, /*DontDefer=*/false,
IsForDefinition);
- } else if (isa<FunctionDecl>(D)) {
+ }
+
+ if (isa<FunctionDecl>(D)) {
const CGFunctionInfo &FI = getTypes().arrangeGlobalDeclaration(GD);
llvm::FunctionType *Ty = getTypes().GetFunctionType(FI);
return GetAddrOfFunction(GD, Ty, /*ForVTable=*/false, /*DontDefer=*/false,
IsForDefinition);
- } else
- return GetAddrOfGlobalVar(cast<VarDecl>(D), /*Ty=*/nullptr,
- IsForDefinition);
+ }
+
+ return GetAddrOfGlobalVar(cast<VarDecl>(D), /*Ty=*/nullptr, IsForDefinition);
}
llvm::GlobalVariable *CodeGenModule::CreateOrReplaceCXXRuntimeVariable(
@@ -3619,7 +4021,7 @@ llvm::GlobalVariable *CodeGenModule::CreateOrReplaceCXXRuntimeVariable(
if (GV) {
// Check if the variable has the right type.
- if (GV->getType()->getElementType() == Ty)
+ if (GV->getValueType() == Ty)
return GV;
// Because C++ name mangling, the only way we can end up with an already
@@ -3668,11 +4070,10 @@ llvm::Constant *CodeGenModule::GetAddrOfGlobalVar(const VarDecl *D,
if (!Ty)
Ty = getTypes().ConvertTypeForMem(ASTTy);
- llvm::PointerType *PTy =
- llvm::PointerType::get(Ty, getContext().getTargetAddressSpace(ASTTy));
-
StringRef MangledName = getMangledName(D);
- return GetOrCreateLLVMGlobal(MangledName, PTy, D, IsForDefinition);
+ return GetOrCreateLLVMGlobal(MangledName, Ty,
+ getContext().getTargetAddressSpace(ASTTy), D,
+ IsForDefinition);
}
/// CreateRuntimeVariable - Create a new runtime global variable with the
@@ -3680,12 +4081,11 @@ llvm::Constant *CodeGenModule::GetAddrOfGlobalVar(const VarDecl *D,
llvm::Constant *
CodeGenModule::CreateRuntimeVariable(llvm::Type *Ty,
StringRef Name) {
- auto PtrTy =
+ auto AddrSpace =
getContext().getLangOpts().OpenCL
- ? llvm::PointerType::get(
- Ty, getContext().getTargetAddressSpace(LangAS::opencl_global))
- : llvm::PointerType::getUnqual(Ty);
- auto *Ret = GetOrCreateLLVMGlobal(Name, PtrTy, nullptr);
+ ? getContext().getTargetAddressSpace(LangAS::opencl_global)
+ : 0;
+ auto *Ret = GetOrCreateLLVMGlobal(Name, Ty, AddrSpace, nullptr);
setDSOLocal(cast<llvm::GlobalValue>(Ret->stripPointerCasts()));
return Ret;
}
@@ -3713,6 +4113,10 @@ void CodeGenModule::EmitTentativeDefinition(const VarDecl *D) {
EmitGlobalVarDefinition(D);
}
+void CodeGenModule::EmitExternalDeclaration(const VarDecl *D) {
+ EmitExternalVarDeclaration(D);
+}
+
CharUnits CodeGenModule::GetTargetTypeStoreSize(llvm::Type *Ty) const {
return Context.toCharUnitsFromBits(
getDataLayout().getTypeStoreSizeInBits(Ty));
@@ -3723,12 +4127,18 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
if (LangOpts.OpenCL) {
AddrSpace = D ? D->getType().getAddressSpace() : LangAS::opencl_global;
assert(AddrSpace == LangAS::opencl_global ||
+ AddrSpace == LangAS::opencl_global_device ||
+ AddrSpace == LangAS::opencl_global_host ||
AddrSpace == LangAS::opencl_constant ||
AddrSpace == LangAS::opencl_local ||
AddrSpace >= LangAS::FirstTargetAddressSpace);
return AddrSpace;
}
+ if (LangOpts.SYCLIsDevice &&
+ (!D || D->getType().getAddressSpace() == LangAS::Default))
+ return LangAS::sycl_global;
+
if (LangOpts.CUDA && LangOpts.CUDAIsDevice) {
if (D && D->hasAttr<CUDAConstantAttr>())
return LangAS::cuda_constant;
@@ -3750,10 +4160,12 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
return getTargetCodeGenInfo().getGlobalVarAddressSpace(*this, D);
}
-LangAS CodeGenModule::getStringLiteralAddressSpace() const {
+LangAS CodeGenModule::GetGlobalConstantAddressSpace() const {
// OpenCL v1.2 s6.5.3: a string literal is in the constant address space.
if (LangOpts.OpenCL)
return LangAS::opencl_constant;
+ if (LangOpts.SYCLIsDevice)
+ return LangAS::sycl_global;
if (auto AS = getTarget().getConstantAddressSpace())
return AS.getValue();
return LangAS::Default;
@@ -3772,13 +4184,12 @@ castStringLiteralToDefaultAddressSpace(CodeGenModule &CGM,
llvm::GlobalVariable *GV) {
llvm::Constant *Cast = GV;
if (!CGM.getLangOpts().OpenCL) {
- if (auto AS = CGM.getTarget().getConstantAddressSpace()) {
- if (AS != LangAS::Default)
- Cast = CGM.getTargetCodeGenInfo().performAddrSpaceCast(
- CGM, GV, AS.getValue(), LangAS::Default,
- GV->getValueType()->getPointerTo(
- CGM.getContext().getTargetAddressSpace(LangAS::Default)));
- }
+ auto AS = CGM.GetGlobalConstantAddressSpace();
+ if (AS != LangAS::Default)
+ Cast = CGM.getTargetCodeGenInfo().performAddrSpaceCast(
+ CGM, GV, AS, LangAS::Default,
+ GV->getValueType()->getPointerTo(
+ CGM.getContext().getTargetAddressSpace(LangAS::Default)));
}
return Cast;
}
@@ -3885,17 +4296,20 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
getLangOpts().CUDAIsDevice && D->hasAttr<CUDASharedAttr>();
// Shadows of initialized device-side global variables are also left
// undefined.
+ // Managed Variables should be initialized on both host side and device side.
bool IsCUDAShadowVar =
- !getLangOpts().CUDAIsDevice &&
+ !getLangOpts().CUDAIsDevice && !D->hasAttr<HIPManagedAttr>() &&
(D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>() ||
D->hasAttr<CUDASharedAttr>());
- // HIP pinned shadow of initialized host-side global variables are also
- // left undefined.
- bool IsHIPPinnedShadowVar =
- getLangOpts().CUDAIsDevice && D->hasAttr<HIPPinnedShadowAttr>();
+ bool IsCUDADeviceShadowVar =
+ getLangOpts().CUDAIsDevice && !D->hasAttr<HIPManagedAttr>() &&
+ (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
+ D->getType()->isCUDADeviceBuiltinTextureType());
if (getLangOpts().CUDA &&
- (IsCUDASharedVar || IsCUDAShadowVar || IsHIPPinnedShadowVar))
- Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
+ (IsCUDASharedVar || IsCUDAShadowVar || IsCUDADeviceShadowVar))
+ Init = llvm::UndefValue::get(getTypes().ConvertTypeForMem(ASTTy));
+ else if (D->hasAttr<LoaderUninitializedAttr>())
+ Init = llvm::UndefValue::get(getTypes().ConvertTypeForMem(ASTTy));
else if (!InitExpr) {
// This is a tentative definition; tentative definitions are
// implicitly initialized with { 0 }.
@@ -3953,7 +4367,7 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
// "extern int x[];") and then a definition of a different type (e.g.
// "int x[10];"). This also happens when an initializer has a different type
// from the type of the global (this happens with unions).
- if (!GV || GV->getType()->getElementType() != InitType ||
+ if (!GV || GV->getValueType() != InitType ||
GV->getType()->getAddressSpace() !=
getContext().getTargetAddressSpace(GetGlobalVarAddressSpace(D))) {
@@ -3967,7 +4381,8 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
// Replace all uses of the old global with the new global
llvm::Constant *NewPtrForOldDecl =
- llvm::ConstantExpr::getBitCast(GV, Entry->getType());
+ llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(GV,
+ Entry->getType());
Entry->replaceAllUsesWith(NewPtrForOldDecl);
// Erase the old global, since it is no longer used.
@@ -3996,38 +4411,14 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
(D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()))
GV->setExternallyInitialized(true);
} else {
- // Host-side shadows of external declarations of device-side
- // global variables become internal definitions. These have to
- // be internal in order to prevent name conflicts with global
- // host variables with the same name in a different TUs.
- if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
- D->hasAttr<HIPPinnedShadowAttr>()) {
- Linkage = llvm::GlobalValue::InternalLinkage;
-
- // Shadow variables and their properties must be registered
- // with CUDA runtime.
- unsigned Flags = 0;
- if (!D->hasDefinition())
- Flags |= CGCUDARuntime::ExternDeviceVar;
- if (D->hasAttr<CUDAConstantAttr>())
- Flags |= CGCUDARuntime::ConstantDeviceVar;
- // Extern global variables will be registered in the TU where they are
- // defined.
- if (!D->hasExternalStorage())
- getCUDARuntime().registerDeviceVar(D, *GV, Flags);
- } else if (D->hasAttr<CUDASharedAttr>())
- // __shared__ variables are odd. Shadows do get created, but
- // they are not registered with the CUDA runtime, so they
- // can't really be used to access their device-side
- // counterparts. It's not clear yet whether it's nvcc's bug or
- // a feature, but we've got to do the same for compatibility.
- Linkage = llvm::GlobalValue::InternalLinkage;
+ getCUDARuntime().internalizeDeviceSideVar(D, Linkage);
}
+ getCUDARuntime().handleVarRegistration(D, *GV);
}
- if (!IsHIPPinnedShadowVar)
- GV->setInitializer(Init);
- if (emitter) emitter->finalize(GV);
+ GV->setInitializer(Init);
+ if (emitter)
+ emitter->finalize(GV);
// If it is safe to mark the global 'constant', do so now.
GV->setConstant(!NeedsGlobalCtor && !NeedsGlobalDtor &&
@@ -4042,17 +4433,24 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
GV->setAlignment(getContext().getDeclAlign(D).getAsAlign());
- // On Darwin, if the normal linkage of a C++ thread_local variable is
- // LinkOnce or Weak, we keep the normal linkage to prevent multiple
- // copies within a linkage unit; otherwise, the backing variable has
- // internal linkage and all accesses should just be calls to the
- // Itanium-specified entry point, which has the normal linkage of the
- // variable. This is to preserve the ability to change the implementation
- // behind the scenes.
- if (!D->isStaticLocal() && D->getTLSKind() == VarDecl::TLS_Dynamic &&
+ // On Darwin, unlike other Itanium C++ ABI platforms, the thread-wrapper
+ // function is only defined alongside the variable, not also alongside
+ // callers. Normally, all accesses to a thread_local go through the
+ // thread-wrapper in order to ensure initialization has occurred, underlying
+ // variable will never be used other than the thread-wrapper, so it can be
+ // converted to internal linkage.
+ //
+ // However, if the variable has the 'constinit' attribute, it _can_ be
+ // referenced directly, without calling the thread-wrapper, so the linkage
+ // must not be changed.
+ //
+ // Additionally, if the variable isn't plain external linkage, e.g. if it's
+ // weak or linkonce, the de-duplication semantics are important to preserve,
+ // so we don't change the linkage.
+ if (D->getTLSKind() == VarDecl::TLS_Dynamic &&
+ Linkage == llvm::GlobalValue::ExternalLinkage &&
Context.getTargetInfo().getTriple().isOSDarwin() &&
- !llvm::GlobalVariable::isLinkOnceLinkage(Linkage) &&
- !llvm::GlobalVariable::isWeakLinkage(Linkage))
+ !D->hasAttr<ConstInitAttr>())
Linkage = llvm::GlobalValue::InternalLinkage;
GV->setLinkage(Linkage);
@@ -4092,10 +4490,22 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
// Emit global variable debug information.
if (CGDebugInfo *DI = getModuleDebugInfo())
- if (getCodeGenOpts().getDebugInfo() >= codegenoptions::LimitedDebugInfo)
+ if (getCodeGenOpts().hasReducedDebugInfo())
DI->EmitGlobalVariable(GV, D);
}
+void CodeGenModule::EmitExternalVarDeclaration(const VarDecl *D) {
+ if (CGDebugInfo *DI = getModuleDebugInfo())
+ if (getCodeGenOpts().hasReducedDebugInfo()) {
+ QualType ASTTy = D->getType();
+ llvm::Type *Ty = getTypes().ConvertTypeForMem(D->getType());
+ llvm::Constant *GV = GetOrCreateLLVMGlobal(
+ D->getName(), Ty, getContext().getTargetAddressSpace(ASTTy), D);
+ DI->EmitExternalVariable(
+ cast<llvm::GlobalVariable>(GV->stripPointerCasts()), D);
+ }
+}
+
static bool isVarDeclStrongDefinition(const ASTContext &Context,
CodeGenModule &CGM, const VarDecl *D,
bool NoCommon) {
@@ -4213,13 +4623,16 @@ llvm::GlobalValue::LinkageTypes CodeGenModule::getLLVMLinkageForDeclarator(
// and must all be equivalent. However, we are not allowed to
// throw away these explicit instantiations.
//
- // We don't currently support CUDA device code spread out across multiple TUs,
+ // CUDA/HIP: For -fno-gpu-rdc case, device code is limited to one TU,
// so say that CUDA templates are either external (for kernels) or internal.
- // This lets llvm perform aggressive inter-procedural optimizations.
+ // This lets llvm perform aggressive inter-procedural optimizations. For
+ // -fgpu-rdc case, device function calls across multiple TU's are allowed,
+ // therefore we need to follow the normal linkage paradigm.
if (Linkage == GVA_StrongODR) {
- if (Context.getLangOpts().AppleKext)
+ if (getLangOpts().AppleKext)
return llvm::Function::ExternalLinkage;
- if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice)
+ if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+ !getLangOpts().GPURelocatableDeviceCode)
return D->hasAttr<CUDAGlobalAttr>() ? llvm::Function::ExternalLinkage
: llvm::Function::InternalLinkage;
return llvm::Function::WeakODRLinkage;
@@ -4259,7 +4672,6 @@ static void replaceUsesOfNonProtoConstant(llvm::Constant *old,
llvm::Type *newRetTy = newFn->getReturnType();
SmallVector<llvm::Value*, 4> newArgs;
- SmallVector<llvm::OperandBundleDef, 1> newBundles;
for (llvm::Value::use_iterator ui = old->use_begin(), ue = old->use_end();
ui != ue; ) {
@@ -4316,6 +4728,7 @@ static void replaceUsesOfNonProtoConstant(llvm::Constant *old,
newArgs.append(callSite->arg_begin(), callSite->arg_begin() + argNo);
// Copy over any operand bundles.
+ SmallVector<llvm::OperandBundleDef, 1> newBundles;
callSite->getOperandBundlesAsDefs(newBundles);
llvm::CallBase *newCall;
@@ -4382,11 +4795,6 @@ void CodeGenModule::HandleCXXStaticMemberVarInstantiation(VarDecl *VD) {
void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
llvm::GlobalValue *GV) {
- // Check if this must be emitted as declare variant.
- if (LangOpts.OpenMP && OpenMPRuntime &&
- OpenMPRuntime->emitDeclareVariant(GD, /*IsForDefinition=*/true))
- return;
-
const auto *D = cast<FunctionDecl>(GD.getDecl());
// Compute the function info and LLVM type.
@@ -4394,7 +4802,7 @@ void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
llvm::FunctionType *Ty = getTypes().GetFunctionType(FI);
// Get or create the prototype for the function.
- if (!GV || (GV->getType()->getElementType() != Ty))
+ if (!GV || (GV->getValueType() != Ty))
GV = cast<llvm::GlobalValue>(GetAddrOfFunction(GD, Ty, /*ForVTable=*/false,
/*DontDefer=*/true,
ForDefinition));
@@ -4415,10 +4823,12 @@ void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
MaybeHandleStaticInExternC(D, Fn);
-
maybeSetTrivialComdat(*D, *Fn);
- CodeGenFunction(*this).GenerateCode(D, Fn, FI);
+ // Set CodeGen attributes that represent floating point environment.
+ setLLVMFunctionFEnvAttributes(D, Fn);
+
+ CodeGenFunction(*this).GenerateCode(GD, Fn, FI);
setNonAliasAttributes(GD, Fn);
SetLLVMFunctionAttributesForDefinition(D, Fn);
@@ -4426,7 +4836,7 @@ void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
if (const ConstructorAttr *CA = D->getAttr<ConstructorAttr>())
AddGlobalCtor(Fn, CA->getPriority());
if (const DestructorAttr *DA = D->getAttr<DestructorAttr>())
- AddGlobalDtor(Fn, DA->getPriority());
+ AddGlobalDtor(Fn, DA->getPriority(), true);
if (D->hasAttr<AnnotateAttr>())
AddGlobalAnnotations(D, Fn);
}
@@ -4462,16 +4872,18 @@ void CodeGenModule::EmitAliasDefinition(GlobalDecl GD) {
/*ForVTable=*/false);
LT = getFunctionLinkage(GD);
} else {
- Aliasee = GetOrCreateLLVMGlobal(AA->getAliasee(),
- llvm::PointerType::getUnqual(DeclTy),
+ Aliasee = GetOrCreateLLVMGlobal(AA->getAliasee(), DeclTy, 0,
/*D=*/nullptr);
- LT = getLLVMLinkageVarDefinition(cast<VarDecl>(GD.getDecl()),
- D->getType().isConstQualified());
+ if (const auto *VD = dyn_cast<VarDecl>(GD.getDecl()))
+ LT = getLLVMLinkageVarDefinition(VD, D->getType().isConstQualified());
+ else
+ LT = getFunctionLinkage(GD);
}
// Create the new alias itself, but don't set a name yet.
+ unsigned AS = Aliasee->getType()->getPointerAddressSpace();
auto *GA =
- llvm::GlobalAlias::create(DeclTy, 0, LT, "", Aliasee, &getModule());
+ llvm::GlobalAlias::create(DeclTy, AS, LT, "", Aliasee, &getModule());
if (Entry) {
if (GA->getAliasee() == Entry) {
@@ -4676,7 +5088,7 @@ CodeGenModule::GetAddrOfConstantCFString(const StringLiteral *Literal) {
DeclContext *DC = TranslationUnitDecl::castToDeclContext(TUDecl);
const VarDecl *VD = nullptr;
- for (const auto &Result : DC->lookup(&II))
+ for (const auto *Result : DC->lookup(&II))
if ((VD = dyn_cast<VarDecl>(Result)))
break;
@@ -4788,6 +5200,8 @@ CodeGenModule::GetAddrOfConstantCFString(const StringLiteral *Literal) {
switch (Triple.getObjectFormat()) {
case llvm::Triple::UnknownObjectFormat:
llvm_unreachable("unknown file format");
+ case llvm::Triple::GOFF:
+ llvm_unreachable("GOFF is not yet implemented");
case llvm::Triple::XCOFF:
llvm_unreachable("XCOFF is not yet implemented");
case llvm::Triple::COFF:
@@ -4886,7 +5300,7 @@ GenerateStringLiteral(llvm::Constant *C, llvm::GlobalValue::LinkageTypes LT,
CodeGenModule &CGM, StringRef GlobalName,
CharUnits Alignment) {
unsigned AddrSpace = CGM.getContext().getTargetAddressSpace(
- CGM.getStringLiteralAddressSpace());
+ CGM.GetGlobalConstantAddressSpace());
llvm::Module &M = CGM.getModule();
// Create a global variable for this string
@@ -5008,13 +5422,26 @@ ConstantAddress CodeGenModule::GetAddrOfGlobalTemporary(
// If we're not materializing a subobject of the temporary, keep the
// cv-qualifiers from the type of the MaterializeTemporaryExpr.
QualType MaterializedType = Init->getType();
- if (Init == E->GetTemporaryExpr())
+ if (Init == E->getSubExpr())
MaterializedType = E->getType();
CharUnits Align = getContext().getTypeAlignInChars(MaterializedType);
- if (llvm::Constant *Slot = MaterializedGlobalTemporaryMap[E])
- return ConstantAddress(Slot, Align);
+ auto InsertResult = MaterializedGlobalTemporaryMap.insert({E, nullptr});
+ if (!InsertResult.second) {
+ // We've seen this before: either we already created it or we're in the
+ // process of doing so.
+ if (!InsertResult.first->second) {
+ // We recursively re-entered this function, probably during emission of
+ // the initializer. Create a placeholder. We'll clean this up in the
+ // outer call, at the end of this function.
+ llvm::Type *Type = getTypes().ConvertTypeForMem(MaterializedType);
+ InsertResult.first->second = new llvm::GlobalVariable(
+ getModule(), Type, false, llvm::GlobalVariable::InternalLinkage,
+ nullptr);
+ }
+ return ConstantAddress(InsertResult.first->second, Align);
+ }
// FIXME: If an externally-visible declaration extends multiple temporaries,
// we need to give each temporary the same name in every translation unit (and
@@ -5031,7 +5458,7 @@ ConstantAddress CodeGenModule::GetAddrOfGlobalTemporary(
// temporary. Note that this might have a different value from the value
// computed by evaluating the initializer if the surrounding constant
// expression modifies the temporary.
- Value = getContext().getMaterializedTemporaryValue(E, false);
+ Value = E->getOrCreateValue(false);
}
// Try evaluating it now, it might have a constant initializer.
@@ -5093,7 +5520,17 @@ ConstantAddress CodeGenModule::GetAddrOfGlobalTemporary(
*this, GV, AddrSpace, LangAS::Default,
Type->getPointerTo(
getContext().getTargetAddressSpace(LangAS::Default)));
- MaterializedGlobalTemporaryMap[E] = CV;
+
+ // Update the map with the new temporary. If we created a placeholder above,
+ // replace it with the new global now.
+ llvm::Constant *&Entry = MaterializedGlobalTemporaryMap[E];
+ if (Entry) {
+ Entry->replaceAllUsesWith(
+ llvm::ConstantExpr::getBitCast(CV, Entry->getType()));
+ llvm::cast<llvm::GlobalVariable>(Entry)->eraseFromParent();
+ }
+ Entry = CV;
+
return ConstantAddress(CV, Align);
}
@@ -5111,11 +5548,12 @@ void CodeGenModule::EmitObjCPropertyImplementations(const
// we want, that just indicates if the decl came from a
// property. What we want to know is if the method is defined in
// this implementation.
- if (!D->getInstanceMethod(PD->getGetterName()))
+ auto *Getter = PID->getGetterMethodDecl();
+ if (!Getter || Getter->isSynthesizedAccessorStub())
CodeGenFunction(*this).GenerateObjCGetter(
- const_cast<ObjCImplementationDecl *>(D), PID);
- if (!PD->isReadOnly() &&
- !D->getInstanceMethod(PD->getSetterName()))
+ const_cast<ObjCImplementationDecl *>(D), PID);
+ auto *Setter = PID->getSetterMethodDecl();
+ if (!PD->isReadOnly() && (!Setter || Setter->isSynthesizedAccessorStub()))
CodeGenFunction(*this).GenerateObjCSetter(
const_cast<ObjCImplementationDecl *>(D), PID);
}
@@ -5152,12 +5590,13 @@ void CodeGenModule::EmitObjCIvarInitializations(ObjCImplementationDecl *D) {
if (needsDestructMethod(D)) {
IdentifierInfo *II = &getContext().Idents.get(".cxx_destruct");
Selector cxxSelector = getContext().Selectors.getSelector(0, &II);
- ObjCMethodDecl *DTORMethod =
- ObjCMethodDecl::Create(getContext(), D->getLocation(), D->getLocation(),
- cxxSelector, getContext().VoidTy, nullptr, D,
- /*isInstance=*/true, /*isVariadic=*/false,
- /*isPropertyAccessor=*/true, /*isImplicitlyDeclared=*/true,
- /*isDefined=*/false, ObjCMethodDecl::Required);
+ ObjCMethodDecl *DTORMethod = ObjCMethodDecl::Create(
+ getContext(), D->getLocation(), D->getLocation(), cxxSelector,
+ getContext().VoidTy, nullptr, D,
+ /*isInstance=*/true, /*isVariadic=*/false,
+ /*isPropertyAccessor=*/true, /*isSynthesizedAccessorStub=*/false,
+ /*isImplicitlyDeclared=*/true,
+ /*isDefined=*/false, ObjCMethodDecl::Required);
D->addInstanceMethod(DTORMethod);
CodeGenFunction(*this).GenerateObjCCtorDtorMethod(D, DTORMethod, false);
D->setHasDestructors(true);
@@ -5172,17 +5611,13 @@ void CodeGenModule::EmitObjCIvarInitializations(ObjCImplementationDecl *D) {
IdentifierInfo *II = &getContext().Idents.get(".cxx_construct");
Selector cxxSelector = getContext().Selectors.getSelector(0, &II);
// The constructor returns 'self'.
- ObjCMethodDecl *CTORMethod = ObjCMethodDecl::Create(getContext(),
- D->getLocation(),
- D->getLocation(),
- cxxSelector,
- getContext().getObjCIdType(),
- nullptr, D, /*isInstance=*/true,
- /*isVariadic=*/false,
- /*isPropertyAccessor=*/true,
- /*isImplicitlyDeclared=*/true,
- /*isDefined=*/false,
- ObjCMethodDecl::Required);
+ ObjCMethodDecl *CTORMethod = ObjCMethodDecl::Create(
+ getContext(), D->getLocation(), D->getLocation(), cxxSelector,
+ getContext().getObjCIdType(), nullptr, D, /*isInstance=*/true,
+ /*isVariadic=*/false,
+ /*isPropertyAccessor=*/true, /*isSynthesizedAccessorStub=*/false,
+ /*isImplicitlyDeclared=*/true,
+ /*isDefined=*/false, ObjCMethodDecl::Required);
D->addInstanceMethod(CTORMethod);
CodeGenFunction(*this).GenerateObjCCtorDtorMethod(D, CTORMethod, true);
D->setHasNonZeroConstructors(true);
@@ -5191,9 +5626,7 @@ void CodeGenModule::EmitObjCIvarInitializations(ObjCImplementationDecl *D) {
// EmitLinkageSpec - Emit all declarations in a linkage spec.
void CodeGenModule::EmitLinkageSpec(const LinkageSpecDecl *LSD) {
if (LSD->getLanguage() != LinkageSpecDecl::lang_c &&
- LSD->getLanguage() != LinkageSpecDecl::lang_cxx &&
- LSD->getLanguage() != LinkageSpecDecl::lang_cxx_11 &&
- LSD->getLanguage() != LinkageSpecDecl::lang_cxx_14) {
+ LSD->getLanguage() != LinkageSpecDecl::lang_cxx) {
ErrorUnsupported(LSD, "linkage spec");
return;
}
@@ -5223,6 +5656,11 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
if (D->isTemplated())
return;
+ // Consteval function shouldn't be emitted.
+ if (auto *FD = dyn_cast<FunctionDecl>(D))
+ if (FD->isConsteval())
+ return;
+
switch (D->getKind()) {
case Decl::CXXConversion:
case Decl::CXXMethod:
@@ -5258,22 +5696,27 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
break;
case Decl::ClassTemplateSpecialization: {
const auto *Spec = cast<ClassTemplateSpecializationDecl>(D);
- if (DebugInfo &&
- Spec->getSpecializationKind() == TSK_ExplicitInstantiationDefinition &&
- Spec->hasDefinition())
- DebugInfo->completeTemplateDefinition(*Spec);
+ if (CGDebugInfo *DI = getModuleDebugInfo())
+ if (Spec->getSpecializationKind() ==
+ TSK_ExplicitInstantiationDefinition &&
+ Spec->hasDefinition())
+ DI->completeTemplateDefinition(*Spec);
} LLVM_FALLTHROUGH;
- case Decl::CXXRecord:
- if (DebugInfo) {
+ case Decl::CXXRecord: {
+ CXXRecordDecl *CRD = cast<CXXRecordDecl>(D);
+ if (CGDebugInfo *DI = getModuleDebugInfo()) {
+ if (CRD->hasDefinition())
+ DI->EmitAndRetainType(getContext().getRecordType(cast<RecordDecl>(D)));
if (auto *ES = D->getASTContext().getExternalSource())
if (ES->hasExternalDefinitions(D) == ExternalASTSource::EK_Never)
- DebugInfo->completeUnusedClass(cast<CXXRecordDecl>(*D));
+ DI->completeUnusedClass(*CRD);
}
// Emit any static data members, they may be definitions.
- for (auto *I : cast<CXXRecordDecl>(D)->decls())
+ for (auto *I : CRD->decls())
if (isa<VarDecl>(I) || isa<CXXRecordDecl>(I))
EmitTopLevelDecl(I);
break;
+ }
// No code generation needed.
case Decl::UsingShadow:
case Decl::ClassTemplate:
@@ -5289,15 +5732,15 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
case Decl::Using: // using X; [C++]
if (CGDebugInfo *DI = getModuleDebugInfo())
DI->EmitUsingDecl(cast<UsingDecl>(*D));
- return;
+ break;
case Decl::NamespaceAlias:
if (CGDebugInfo *DI = getModuleDebugInfo())
DI->EmitNamespaceAlias(cast<NamespaceAliasDecl>(*D));
- return;
+ break;
case Decl::UsingDirective: // using namespace X; [C++]
if (CGDebugInfo *DI = getModuleDebugInfo())
DI->EmitUsingDirective(cast<UsingDirectiveDecl>(*D));
- return;
+ break;
case Decl::CXXConstructor:
getCXXABI().EmitCXXConstructors(cast<CXXConstructorDecl>(D));
break;
@@ -5336,7 +5779,7 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
ObjCRuntime->GenerateClass(OMD);
// Emit global variable debug information.
if (CGDebugInfo *DI = getModuleDebugInfo())
- if (getCodeGenOpts().getDebugInfo() >= codegenoptions::LimitedDebugInfo)
+ if (getCodeGenOpts().hasReducedDebugInfo())
DI->getOrCreateInterfaceType(getContext().getObjCInterfaceType(
OMD->getClassInterface()), OMD->getLocation());
break;
@@ -5388,6 +5831,9 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
// File-scope asm is ignored during device-side OpenMP compilation.
if (LangOpts.OpenMPIsDevice)
break;
+ // File-scope asm is ignored during device-side SYCL compilation.
+ if (LangOpts.SYCLIsDevice)
+ break;
auto *AD = cast<FileScopeAsmDecl>(D);
getModule().appendModuleInlineAsm(AD->getAsmString()->getString());
break;
@@ -5445,6 +5891,7 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
break;
case Decl::OMPAllocate:
+ EmitOMPAllocateDecl(cast<OMPAllocateDecl>(D));
break;
case Decl::OMPDeclareReduction:
@@ -5459,6 +5906,25 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
EmitOMPRequiresDecl(cast<OMPRequiresDecl>(D));
break;
+ case Decl::Typedef:
+ case Decl::TypeAlias: // using foo = bar; [C++11]
+ if (CGDebugInfo *DI = getModuleDebugInfo())
+ DI->EmitAndRetainType(
+ getContext().getTypedefType(cast<TypedefNameDecl>(D)));
+ break;
+
+ case Decl::Record:
+ if (CGDebugInfo *DI = getModuleDebugInfo())
+ if (cast<RecordDecl>(D)->getDefinition())
+ DI->EmitAndRetainType(getContext().getRecordType(cast<RecordDecl>(D)));
+ break;
+
+ case Decl::Enum:
+ if (CGDebugInfo *DI = getModuleDebugInfo())
+ if (cast<EnumDecl>(D)->getDefinition())
+ DI->EmitAndRetainType(getContext().getEnumType(cast<EnumDecl>(D)));
+ break;
+
default:
// Make sure we handled everything we should, every other kind is a
// non-top-level decl. FIXME: Would be nice to have an isTopLevelDeclKind
@@ -5480,10 +5946,10 @@ void CodeGenModule::AddDeferredUnusedCoverageMapping(Decl *D) {
case Decl::CXXConstructor:
case Decl::CXXDestructor: {
if (!cast<FunctionDecl>(D)->doesThisDeclarationHaveABody())
- return;
+ break;
SourceManager &SM = getContext().getSourceManager();
if (LimitedCoverage && SM.getMainFileID() != SM.getFileID(D->getBeginLoc()))
- return;
+ break;
auto I = DeferredEmptyCoverageMappingDecls.find(D);
if (I == DeferredEmptyCoverageMappingDecls.end())
DeferredEmptyCoverageMappingDecls[D] = true;
@@ -5549,6 +6015,17 @@ void CodeGenModule::EmitDeferredUnusedCoverageMappings() {
}
}
+void CodeGenModule::EmitMainVoidAlias() {
+ // In order to transition away from "__original_main" gracefully, emit an
+ // alias for "main" in the no-argument case so that libc can detect when
+ // new-style no-argument main is in used.
+ if (llvm::Function *F = getModule().getFunction("main")) {
+ if (!F->isDeclaration() && F->arg_size() == 0 && !F->isVarArg() &&
+ F->getReturnType()->isIntegerTy(Context.getTargetInfo().getIntWidth()))
+ addUsedGlobal(llvm::GlobalAlias::create("__main_void", F));
+ }
+}
+
/// Turns the given pointer into a constant.
static llvm::Constant *GetPointerConstant(llvm::LLVMContext &Context,
const void *Ptr) {
@@ -5584,7 +6061,7 @@ void CodeGenModule::EmitStaticExternCAliases() {
IdentifierInfo *Name = I.first;
llvm::GlobalValue *Val = I.second;
if (Val && !getModule().getNamedValue(Name->getName()))
- addUsedGlobal(llvm::GlobalAlias::create(Name->getName(), Val));
+ addCompilerUsedGlobal(llvm::GlobalAlias::create(Name->getName(), Val));
}
}
@@ -5663,21 +6140,6 @@ void CodeGenModule::EmitCommandLineMetadata() {
CommandLineMetadata->addOperand(llvm::MDNode::get(Ctx, CommandLineNode));
}
-void CodeGenModule::EmitTargetMetadata() {
- // Warning, new MangledDeclNames may be appended within this loop.
- // We rely on MapVector insertions adding new elements to the end
- // of the container.
- // FIXME: Move this loop into the one target that needs it, and only
- // loop over those declarations for which we couldn't emit the target
- // metadata when we emitted the declaration.
- for (unsigned I = 0; I != MangledDeclNames.size(); ++I) {
- auto Val = *(MangledDeclNames.begin() + I);
- const Decl *D = Val.first.getDecl()->getMostRecentDecl();
- llvm::GlobalValue *GV = GetGlobalValue(Val.second);
- getTargetCodeGenInfo().emitTargetMD(D, GV, *this);
- }
-}
-
void CodeGenModule::EmitCoverageFile() {
if (getCodeGenOpts().CoverageDataFile.empty() &&
getCodeGenOpts().CoverageNotesFile.empty())
@@ -5700,39 +6162,14 @@ void CodeGenModule::EmitCoverageFile() {
}
}
-llvm::Constant *CodeGenModule::EmitUuidofInitializer(StringRef Uuid) {
- // Sema has checked that all uuid strings are of the form
- // "12345678-1234-1234-1234-1234567890ab".
- assert(Uuid.size() == 36);
- for (unsigned i = 0; i < 36; ++i) {
- if (i == 8 || i == 13 || i == 18 || i == 23) assert(Uuid[i] == '-');
- else assert(isHexDigit(Uuid[i]));
- }
-
- // The starts of all bytes of Field3 in Uuid. Field 3 is "1234-1234567890ab".
- const unsigned Field3ValueOffsets[8] = { 19, 21, 24, 26, 28, 30, 32, 34 };
-
- llvm::Constant *Field3[8];
- for (unsigned Idx = 0; Idx < 8; ++Idx)
- Field3[Idx] = llvm::ConstantInt::get(
- Int8Ty, Uuid.substr(Field3ValueOffsets[Idx], 2), 16);
-
- llvm::Constant *Fields[4] = {
- llvm::ConstantInt::get(Int32Ty, Uuid.substr(0, 8), 16),
- llvm::ConstantInt::get(Int16Ty, Uuid.substr(9, 4), 16),
- llvm::ConstantInt::get(Int16Ty, Uuid.substr(14, 4), 16),
- llvm::ConstantArray::get(llvm::ArrayType::get(Int8Ty, 8), Field3)
- };
-
- return llvm::ConstantStruct::getAnon(Fields);
-}
-
llvm::Constant *CodeGenModule::GetAddrOfRTTIDescriptor(QualType Ty,
bool ForEH) {
// Return a bogus pointer if RTTI is disabled, unless it's for EH.
// FIXME: should we even be calling this method if RTTI is disabled
// and it's not for EH?
- if ((!ForEH && !getLangOpts().RTTI) || getLangOpts().CUDAIsDevice)
+ if ((!ForEH && !getLangOpts().RTTI) || getLangOpts().CUDAIsDevice ||
+ (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
+ getTriple().isNVPTX()))
return llvm::Constant::getNullValue(Int8PtrTy);
if (ForEH && Ty->isObjCObjectPointerType() &&
@@ -5860,71 +6297,122 @@ void CodeGenModule::AddVTableTypeMetadata(llvm::GlobalVariable *VTable,
}
}
-TargetAttr::ParsedTargetAttr CodeGenModule::filterFunctionTargetAttrs(const TargetAttr *TD) {
- assert(TD != nullptr);
- TargetAttr::ParsedTargetAttr ParsedAttr = TD->parse();
-
- ParsedAttr.Features.erase(
- llvm::remove_if(ParsedAttr.Features,
- [&](const std::string &Feat) {
- return !Target.isValidFeatureName(
- StringRef{Feat}.substr(1));
- }),
- ParsedAttr.Features.end());
- return ParsedAttr;
-}
-
-
-// Fills in the supplied string map with the set of target features for the
-// passed in function.
-void CodeGenModule::getFunctionFeatureMap(llvm::StringMap<bool> &FeatureMap,
- GlobalDecl GD) {
- StringRef TargetCPU = Target.getTargetOpts().CPU;
- const FunctionDecl *FD = GD.getDecl()->getAsFunction();
- if (const auto *TD = FD->getAttr<TargetAttr>()) {
- TargetAttr::ParsedTargetAttr ParsedAttr = filterFunctionTargetAttrs(TD);
-
- // Make a copy of the features as passed on the command line into the
- // beginning of the additional features from the function to override.
- ParsedAttr.Features.insert(ParsedAttr.Features.begin(),
- Target.getTargetOpts().FeaturesAsWritten.begin(),
- Target.getTargetOpts().FeaturesAsWritten.end());
-
- if (ParsedAttr.Architecture != "" &&
- Target.isValidCPUName(ParsedAttr.Architecture))
- TargetCPU = ParsedAttr.Architecture;
-
- // Now populate the feature map, first with the TargetCPU which is either
- // the default or a new one from the target attribute string. Then we'll use
- // the passed in features (FeaturesAsWritten) along with the new ones from
- // the attribute.
- Target.initFeatureMap(FeatureMap, getDiags(), TargetCPU,
- ParsedAttr.Features);
- } else if (const auto *SD = FD->getAttr<CPUSpecificAttr>()) {
- llvm::SmallVector<StringRef, 32> FeaturesTmp;
- Target.getCPUSpecificCPUDispatchFeatures(
- SD->getCPUName(GD.getMultiVersionIndex())->getName(), FeaturesTmp);
- std::vector<std::string> Features(FeaturesTmp.begin(), FeaturesTmp.end());
- Target.initFeatureMap(FeatureMap, getDiags(), TargetCPU, Features);
- } else {
- Target.initFeatureMap(FeatureMap, getDiags(), TargetCPU,
- Target.getTargetOpts().Features);
- }
-}
-
llvm::SanitizerStatReport &CodeGenModule::getSanStats() {
if (!SanStats)
SanStats = std::make_unique<llvm::SanitizerStatReport>(&getModule());
return *SanStats;
}
+
llvm::Value *
CodeGenModule::createOpenCLIntToSamplerConversion(const Expr *E,
CodeGenFunction &CGF) {
llvm::Constant *C = ConstantEmitter(CGF).emitAbstract(E, E->getType());
- auto SamplerT = getOpenCLRuntime().getSamplerType(E->getType().getTypePtr());
- auto FTy = llvm::FunctionType::get(SamplerT, {C->getType()}, false);
- return CGF.Builder.CreateCall(CreateRuntimeFunction(FTy,
- "__translate_sampler_initializer"),
- {C});
+ auto *SamplerT = getOpenCLRuntime().getSamplerType(E->getType().getTypePtr());
+ auto *FTy = llvm::FunctionType::get(SamplerT, {C->getType()}, false);
+ auto *Call = CGF.EmitRuntimeCall(
+ CreateRuntimeFunction(FTy, "__translate_sampler_initializer"), {C});
+ return Call;
+}
+
+CharUnits CodeGenModule::getNaturalPointeeTypeAlignment(
+ QualType T, LValueBaseInfo *BaseInfo, TBAAAccessInfo *TBAAInfo) {
+ return getNaturalTypeAlignment(T->getPointeeType(), BaseInfo, TBAAInfo,
+ /* forPointeeType= */ true);
+}
+
+CharUnits CodeGenModule::getNaturalTypeAlignment(QualType T,
+ LValueBaseInfo *BaseInfo,
+ TBAAAccessInfo *TBAAInfo,
+ bool forPointeeType) {
+ if (TBAAInfo)
+ *TBAAInfo = getTBAAAccessInfo(T);
+
+ // FIXME: This duplicates logic in ASTContext::getTypeAlignIfKnown. But
+ // that doesn't return the information we need to compute BaseInfo.
+
+ // Honor alignment typedef attributes even on incomplete types.
+ // We also honor them straight for C++ class types, even as pointees;
+ // there's an expressivity gap here.
+ if (auto TT = T->getAs<TypedefType>()) {
+ if (auto Align = TT->getDecl()->getMaxAlignment()) {
+ if (BaseInfo)
+ *BaseInfo = LValueBaseInfo(AlignmentSource::AttributedType);
+ return getContext().toCharUnitsFromBits(Align);
+ }
+ }
+
+ bool AlignForArray = T->isArrayType();
+
+ // Analyze the base element type, so we don't get confused by incomplete
+ // array types.
+ T = getContext().getBaseElementType(T);
+
+ if (T->isIncompleteType()) {
+ // We could try to replicate the logic from
+ // ASTContext::getTypeAlignIfKnown, but nothing uses the alignment if the
+ // type is incomplete, so it's impossible to test. We could try to reuse
+ // getTypeAlignIfKnown, but that doesn't return the information we need
+ // to set BaseInfo. So just ignore the possibility that the alignment is
+ // greater than one.
+ if (BaseInfo)
+ *BaseInfo = LValueBaseInfo(AlignmentSource::Type);
+ return CharUnits::One();
+ }
+
+ if (BaseInfo)
+ *BaseInfo = LValueBaseInfo(AlignmentSource::Type);
+
+ CharUnits Alignment;
+ const CXXRecordDecl *RD;
+ if (T.getQualifiers().hasUnaligned()) {
+ Alignment = CharUnits::One();
+ } else if (forPointeeType && !AlignForArray &&
+ (RD = T->getAsCXXRecordDecl())) {
+ // For C++ class pointees, we don't know whether we're pointing at a
+ // base or a complete object, so we generally need to use the
+ // non-virtual alignment.
+ Alignment = getClassPointerAlignment(RD);
+ } else {
+ Alignment = getContext().getTypeAlignInChars(T);
+ }
+
+ // Cap to the global maximum type alignment unless the alignment
+ // was somehow explicit on the type.
+ if (unsigned MaxAlign = getLangOpts().MaxTypeAlign) {
+ if (Alignment.getQuantity() > MaxAlign &&
+ !getContext().isAlignmentRequired(T))
+ Alignment = CharUnits::fromQuantity(MaxAlign);
+ }
+ return Alignment;
+}
+
+bool CodeGenModule::stopAutoInit() {
+ unsigned StopAfter = getContext().getLangOpts().TrivialAutoVarInitStopAfter;
+ if (StopAfter) {
+ // This number is positive only when -ftrivial-auto-var-init-stop-after=* is
+ // used
+ if (NumAutoVarInit >= StopAfter) {
+ return true;
+ }
+ if (!NumAutoVarInit) {
+ unsigned DiagID = getDiags().getCustomDiagID(
+ DiagnosticsEngine::Warning,
+ "-ftrivial-auto-var-init-stop-after=%0 has been enabled to limit the "
+ "number of times ftrivial-auto-var-init=%1 gets applied.");
+ getDiags().Report(DiagID)
+ << StopAfter
+ << (getContext().getLangOpts().getTrivialAutoVarInit() ==
+ LangOptions::TrivialAutoVarInitKind::Zero
+ ? "zero"
+ : "pattern");
+ }
+ ++NumAutoVarInit;
+ }
+ return false;
+}
+
+void CodeGenModule::printPostfixForExternalizedStaticVar(
+ llvm::raw_ostream &OS) const {
+ OS << ".static." << getContext().getCUIDHash();
}