From 13fbcb42fff2606c3ce926a7c4bb8989aa18758c Mon Sep 17 00:00:00 2001 From: joerg Date: Sun, 30 May 2021 01:25:15 +0000 Subject: Import clang 249b40b558955afe5ac2b549edcf2d7f859c8cc9. --- .../llvm/dist/clang/lib/CodeGen/CodeGenModule.cpp | 1576 +++++++++++++------- 1 file changed, 1032 insertions(+), 544 deletions(-) (limited to 'external/apache2/llvm/dist/clang/lib/CodeGen/CodeGenModule.cpp') 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 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(Entry); + auto *Alias = cast(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(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(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(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(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(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()) { @@ -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(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(ND)) - MC.mangleCXXCtor(D, GD.getCtorType(), Out); - else if (const auto *D = dyn_cast(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(ND); if (FD && FD->getType()->castAs()->getCallConv() == CC_X86RegCall) { - llvm::raw_svector_ostream Out(Buffer); Out << "__regcall3__" << II->getName(); + } else if (FD && FD->hasAttr() && + 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(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(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(GD.getDecl())) - if (!getLangOpts().CUDAIsDevice && FD->hasAttr()) - 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(ND) || !ND->hasAttr() || + 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(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(ty)) + PDecl = TD->getDecl(); + const OpenCLAccessAttr *A = PDecl->getAttr(); + 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() - ->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() - ->getElementType() - .getCanonicalType() - .getAsString(Policy); - else - baseTypeName = - ty.getUnqualifiedType().getCanonicalType().getAsString(Policy); + ty = isPipe ? ty->castAs()->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(ty)) - PDecl = TD->getDecl(); - const OpenCLAccessAttr *A = PDecl->getAttr(); - 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(); - ShouldAddOptNone &= !F->hasFnAttribute(llvm::Attribute::AlwaysInline); ShouldAddOptNone &= !D->hasAttr(); - if (ShouldAddOptNone || D->hasAttr()) { + // Add optnone, but do so only if the function isn't always_inline. + if ((ShouldAddOptNone || D->hasAttr()) && + !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()) { B.addAttribute(llvm::Attribute::NoDuplicate); - } else if (D->hasAttr()) { + } else if (D->hasAttr() && !F->hasFnAttribute(llvm::Attribute::AlwaysInline)) { + // Add noinline if the function isn't always_inline. B.addAttribute(llvm::Attribute::NoInline); } else if (D->hasAttr() && !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()) + B.addAttribute(llvm::Attribute::Hot); if (D->hasAttr()) 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()) { + 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(D)) @@ -1643,13 +1900,13 @@ void CodeGenModule::SetCommonAttributes(GlobalDecl GD, llvm::GlobalValue *GV) { GV->setVisibility(llvm::GlobalValue::DefaultVisibility); if (D && D->hasAttr()) - addUsedGlobal(GV); + addUsedOrCompilerUsedGlobal(GV); if (CodeGenOpts.KeepStaticConsts && D && isa(D)) { const auto *VD = cast(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 Features; const auto *FD = dyn_cast_or_null(GD.getDecl()); FD = FD ? FD->getMostRecentDecl() : FD; @@ -1667,7 +1925,7 @@ bool CodeGenModule::GetCPUAndFeaturesAttributes(GlobalDecl GD, bool AddedAttr = false; if (TD || SD) { llvm::StringMap FeatureMap; - getFunctionFeatureMap(FeatureMap, GD); + getContext().getFunctionFeatureMap(FeatureMap, GD); // Produce the canonical string for this set of features. for (const llvm::StringMap::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(GO)) { + if (D->hasAttr()) + addUsedGlobal(GV); if (auto *SA = D->getAttr()) GV->addAttribute("bss-section", SA->getName()); if (auto *SA = D->getAttr()) @@ -1720,6 +1989,8 @@ void CodeGenModule::setNonAliasAttributes(GlobalDecl GD, } if (auto *F = dyn_cast(GO)) { + if (D->hasAttr()) + addUsedGlobal(F); if (auto *SA = D->getAttr()) if (!D->getAttr()) 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(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()) 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(FD) || isa(FD)) @@ -1883,7 +2161,7 @@ void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F, } void CodeGenModule::addUsedGlobal(llvm::GlobalValue *GV) { - assert(!GV->isDeclaration() && + assert((isa(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(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 &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 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(E)->getAPValueResult()); + } + llvm::Constant *&Lookup = AnnotationArgs[ID.ComputeHash()]; + if (Lookup) + return Lookup; + + llvm::SmallVector LLVMArgs; + LLVMArgs.reserve(Exprs.size()); + ConstantEmitter ConstEmiter(*this); + llvm::transform(Exprs, std::back_inserter(LLVMArgs), [&](const Expr *E) { + const auto *CE = cast(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(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 V = ProfileList.isFunctionExcluded(Fn->getName(), Kind); + if (V.hasValue()) + return *V; + // Next, check the source location. + if (Loc.isValid()) { + Optional 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 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 ActiveAttr = + OMPDeclareTargetDeclAttr::getActiveAttr(Global); + if (!ActiveAttr || (*ActiveAttr)->getLevel() != (unsigned)-1) + return false; + } + if (const auto *FD = dyn_cast(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(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(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(VD)), /*ForVTable=*/false); else - Aliasee = GetOrCreateLLVMGlobal(AA->getAliasee(), - llvm::PointerType::getUnqual(DeclTy), - nullptr); + Aliasee = GetOrCreateLLVMGlobal(AA->getAliasee(), DeclTy, 0, nullptr); auto *F = cast(Aliasee); F->setLinkage(llvm::Function::ExternalWeakLinkage); @@ -2439,7 +2859,8 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) { !Global->hasAttr() && !Global->hasAttr() && !Global->hasAttr() && - !(LangOpts.HIP && Global->hasAttr())) + !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(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(Global) && @@ -2697,7 +3113,7 @@ bool CodeGenModule::shouldEmitFunction(GlobalDecl GD) { if (CodeGenOpts.OptimizationLevel == 0 && !F->hasAttr()) return false; - if (F->hasAttr()) { + if (F->hasAttr() && !F->hasAttr()) { // Check whether it would be safe to inline this dllimport function. DLLImportFunctionVisitor Visitor; Visitor.TraverseFunctionDecl(const_cast(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(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(GetOrCreateLLVMFunction( - getMangledName(OldGD), Ty, GlobalDecl(), /*ForVTable=*/false, - /*DontDefer=*/true, /*IsThunk=*/false, llvm::AttributeList(), - ForDefinition)); - SetFunctionAttributes(OldGD, cast(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(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()) - AddGlobalAnnotations(D, Fn); -} - void CodeGenModule::EmitGlobalDefinition(GlobalDecl GD, llvm::GlobalValue *GV) { const auto *D = cast(GD.getDecl()); @@ -2851,7 +3223,9 @@ TargetMVPriority(const TargetInfo &TI, } void CodeGenModule::emitMultiVersionFunctions() { - for (GlobalDecl GD : MultiVersionFuncs) { + std::vector MVFuncsToEmit; + MultiVersionFuncs.swap(MVFuncsToEmit); + for (GlobalDecl GD : MVFuncsToEmit) { SmallVector Options; const FunctionDecl *FD = cast(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(); - if (TA && TA->isDefaultVersion()) + if (FD->hasAttr()) UpdateMultiVersionNames(GD, FD); if (!IsForDefinition) return GetOrCreateMultiVersionResolver(GD, Ty, FD); @@ -3147,7 +3527,7 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction( } if ((isa(Entry) || isa(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(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(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(GD.getDecl())->hasAttr()) { + auto *Handle = getCUDARuntime().getKernelHandle( + cast(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(Result)) + for (const auto *Result : DC->lookup(&CII)) + if (const auto *FD = dyn_cast(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(Result); - if (auto LSD = dyn_cast(Result)) - for (const auto &Result : LSD->lookup(&NS)) + for (const auto *Result : DC->lookup(&NS)) { + const NamespaceDecl *ND = dyn_cast(Result); + if (auto *LSD = dyn_cast(Result)) + for (const auto *Result : LSD->lookup(&NS)) if ((ND = dyn_cast(Result))) break; if (ND) - for (const auto &Result : ND->lookup(&CXXII)) + for (const auto *Result : ND->lookup(&CXXII)) if (const auto *FD = dyn_cast(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() && + 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(D) || isa(D)) return getAddrOfCXXStructor(GD, /*FnInfo=*/nullptr, /*FnType=*/nullptr, /*DontDefer=*/false, IsForDefinition); - else if (isa(D)) { - auto FInfo = &getTypes().arrangeCXXMethodDeclaration( - cast(D)); + + if (isa(D)) { + auto FInfo = + &getTypes().arrangeCXXMethodDeclaration(cast(D)); auto Ty = getTypes().GetFunctionType(*FInfo); return GetAddrOfFunction(GD, Ty, /*ForVTable=*/false, /*DontDefer=*/false, IsForDefinition); - } else if (isa(D)) { + } + + if (isa(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(D), /*Ty=*/nullptr, - IsForDefinition); + } + + return GetAddrOfGlobalVar(cast(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(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()) 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(); // 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() && (D->hasAttr() || D->hasAttr() || D->hasAttr()); - // HIP pinned shadow of initialized host-side global variables are also - // left undefined. - bool IsHIPPinnedShadowVar = - getLangOpts().CUDAIsDevice && D->hasAttr(); + bool IsCUDADeviceShadowVar = + getLangOpts().CUDAIsDevice && !D->hasAttr() && + (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()) + 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() || D->hasAttr())) 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() || D->hasAttr() || - D->hasAttr()) { - 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()) - 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()) - // __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()) 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(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() ? 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 newArgs; - SmallVector 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 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(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(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()) AddGlobalCtor(Fn, CA->getPriority()); if (const DestructorAttr *DA = D->getAttr()) - AddGlobalDtor(Fn, DA->getPriority()); + AddGlobalDtor(Fn, DA->getPriority(), true); if (D->hasAttr()) 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(GD.getDecl()), - D->getType().isConstQualified()); + if (const auto *VD = dyn_cast(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(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(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(D), PID); - if (!PD->isReadOnly() && - !D->getInstanceMethod(PD->getSetterName())) + const_cast(D), PID); + auto *Setter = PID->getSetterMethodDecl(); + if (!PD->isReadOnly() && (!Setter || Setter->isSynthesizedAccessorStub())) CodeGenFunction(*this).GenerateObjCSetter( const_cast(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(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(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(D); + if (CGDebugInfo *DI = getModuleDebugInfo()) { + if (CRD->hasDefinition()) + DI->EmitAndRetainType(getContext().getRecordType(cast(D))); if (auto *ES = D->getASTContext().getExternalSource()) if (ES->hasExternalDefinitions(D) == ExternalASTSource::EK_Never) - DebugInfo->completeUnusedClass(cast(*D)); + DI->completeUnusedClass(*CRD); } // Emit any static data members, they may be definitions. - for (auto *I : cast(D)->decls()) + for (auto *I : CRD->decls()) if (isa(I) || isa(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(*D)); - return; + break; case Decl::NamespaceAlias: if (CGDebugInfo *DI = getModuleDebugInfo()) DI->EmitNamespaceAlias(cast(*D)); - return; + break; case Decl::UsingDirective: // using namespace X; [C++] if (CGDebugInfo *DI = getModuleDebugInfo()) DI->EmitUsingDirective(cast(*D)); - return; + break; case Decl::CXXConstructor: getCXXABI().EmitCXXConstructors(cast(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(D); getModule().appendModuleInlineAsm(AD->getAsmString()->getString()); break; @@ -5445,6 +5891,7 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) { break; case Decl::OMPAllocate: + EmitOMPAllocateDecl(cast(D)); break; case Decl::OMPDeclareReduction: @@ -5459,6 +5906,25 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) { EmitOMPRequiresDecl(cast(D)); break; + case Decl::Typedef: + case Decl::TypeAlias: // using foo = bar; [C++11] + if (CGDebugInfo *DI = getModuleDebugInfo()) + DI->EmitAndRetainType( + getContext().getTypedefType(cast(D))); + break; + + case Decl::Record: + if (CGDebugInfo *DI = getModuleDebugInfo()) + if (cast(D)->getDefinition()) + DI->EmitAndRetainType(getContext().getRecordType(cast(D))); + break; + + case Decl::Enum: + if (CGDebugInfo *DI = getModuleDebugInfo()) + if (cast(D)->getDefinition()) + DI->EmitAndRetainType(getContext().getEnumType(cast(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(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 &FeatureMap, - GlobalDecl GD) { - StringRef TargetCPU = Target.getTargetOpts().CPU; - const FunctionDecl *FD = GD.getDecl()->getAsFunction(); - if (const auto *TD = FD->getAttr()) { - 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()) { - llvm::SmallVector FeaturesTmp; - Target.getCPUSpecificCPUDispatchFeatures( - SD->getCPUName(GD.getMultiVersionIndex())->getName(), FeaturesTmp); - std::vector 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(&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()) { + 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(); } -- cgit