Import LLVM 10.0.0 release including clang, lld and lldb.

ok hackroom
tested by plenty

1 files changed, 469 insertions, 1001 deletions

diff --git a/gnu/llvm/clang/lib/CodeGen/CodeGenModule.cpp b/gnu/llvm/clang/lib/CodeGen/CodeGenModule.cpp
index 49a1396b58e..f8866ac4f7f 100644
--- a/gnu/llvm/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/gnu/llvm/clang/lib/CodeGen/CodeGenModule.cpp
@@ -19,7 +19,6 @@
 #include "CGObjCRuntime.h"
 #include "CGOpenCLRuntime.h"
 #include "CGOpenMPRuntime.h"
-#include "CGOpenMPRuntimeAMDGCN.h"
 #include "CGOpenMPRuntimeNVPTX.h"
 #include "CodeGenFunction.h"
 #include "CodeGenPGO.h"
@@ -39,7 +38,6 @@
 #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"
@@ -75,17 +73,16 @@ static llvm::cl::opt<bool> LimitedCoverage(
 static const char AnnotationSection[] = "llvm.metadata";
 
 static CGCXXABI *createCXXABI(CodeGenModule &CGM) {
-  switch (CGM.getContext().getCXXABIKind()) {
-  case TargetCXXABI::AppleARM64:
+  switch (CGM.getTarget().getCXXABI().getKind()) {
   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);
@@ -113,7 +110,6 @@ 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);
@@ -123,8 +119,6 @@ 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());
@@ -180,34 +174,6 @@ 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 : LangOpts.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(toString(IntHash, /* Radix = */ 10, /* Signed = */false))).str();
-  }
 }
 
 CodeGenModule::~CodeGenModule() {}
@@ -246,11 +212,6 @@ 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));
@@ -258,6 +219,14 @@ void CodeGenModule::createOpenMPRuntime() {
       OpenMPRuntime.reset(new CGOpenMPRuntime(*this));
     break;
   }
+
+  // The OpenMP-IR-Builder should eventually replace the above runtime codegens
+  // but we are not there yet so they both reside in CGModule for now and the
+  // OpenMP-IR-Builder is opt-in only.
+  if (LangOpts.OpenMPIRBuilder) {
+    OMPBuilder.reset(new llvm::OpenMPIRBuilder(TheModule));
+    OMPBuilder->initialize();
+  }
 }
 
 void CodeGenModule::createCUDARuntime() {
@@ -402,7 +371,7 @@ void CodeGenModule::checkAliases() {
   for (const GlobalDecl &GD : Aliases) {
     StringRef MangledName = getMangledName(GD);
     llvm::GlobalValue *Entry = GetGlobalValue(MangledName);
-    auto *Alias = cast<llvm::GlobalIndirectSymbol>(Entry);
+    auto *Alias = dyn_cast<llvm::GlobalIndirectSymbol>(Entry);
     Alias->replaceAllUsesWith(llvm::UndefValue::get(Alias->getType()));
     Alias->eraseFromParent();
   }
@@ -431,48 +400,6 @@ 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();
@@ -481,14 +408,16 @@ void CodeGenModule::Release() {
   checkAliases();
   emitMultiVersionFunctions();
   EmitCXXGlobalInitFunc();
-  EmitCXXGlobalCleanUpFunc();
+  EmitCXXGlobalDtorFunc();
   registerGlobalDtorsWithAtExit();
   EmitCXXThreadLocalInitFunc();
   if (ObjCRuntime)
     if (llvm::Function *ObjCInitFunction = ObjCRuntime->ModuleInitFunction())
       AddGlobalCtor(ObjCInitFunction);
-  if (Context.getLangOpts().CUDA && CUDARuntime) {
-    if (llvm::Function *CudaCtorFunction = CUDARuntime->finalizeModule())
+  if (Context.getLangOpts().CUDA && !Context.getLangOpts().CUDAIsDevice &&
+      CUDARuntime) {
+    if (llvm::Function *CudaCtorFunction =
+            CUDARuntime->makeModuleCtorFunction())
       AddGlobalCtor(CudaCtorFunction);
   }
   if (OpenMPRuntime) {
@@ -511,7 +440,6 @@ void CodeGenModule::Release() {
   EmitGlobalAnnotations();
   EmitStaticExternCAliases();
   EmitDeferredUnusedCoverageMappings();
-  CodeGenPGO(*this).setValueProfilingFlag(getModule());
   if (CoverageMapping)
     CoverageMapping->emit();
   if (CodeGenOpts.SanitizeCfiCrossDso) {
@@ -519,26 +447,6 @@ void CodeGenModule::Release() {
     CodeGenFunction(*this).EmitCfiCheckStub();
   }
   emitAtAvailableLinkGuard();
-  if (Context.getTargetInfo().getTriple().isWasm() &&
-      !Context.getTargetInfo().getTriple().isOSEmscripten()) {
-    EmitMainVoidAlias();
-  }
-
-  // Emit reference of __amdgpu_device_library_preserve_asan_functions to
-  // preserve ASAN functions in bitcode libraries.
-  if (LangOpts.Sanitize.has(SanitizerKind::Address) && getTriple().isAMDGPU()) {
-    auto *FT = llvm::FunctionType::get(VoidTy, {});
-    auto *F = llvm::Function::Create(
-        FT, llvm::GlobalValue::ExternalLinkage,
-        "__amdgpu_device_library_preserve_asan_functions", &getModule());
-    auto *Var = new llvm::GlobalVariable(
-        getModule(), FT->getPointerTo(),
-        /*isConstant=*/true, llvm::GlobalValue::WeakAnyLinkage, F,
-        "__amdgpu_device_library_preserve_asan_functions_ptr", nullptr,
-        llvm::GlobalVariable::NotThreadLocal);
-    addCompilerUsedGlobal(Var);
-  }
-
   emitLLVMUsed();
   if (SanStats)
     SanStats->finish();
@@ -575,14 +483,6 @@ void CodeGenModule::Release() {
     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);
@@ -597,10 +497,6 @@ 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
@@ -617,7 +513,7 @@ void CodeGenModule::Release() {
                               "StrictVTablePointersRequirement",
                               llvm::MDNode::get(VMContext, Ops));
   }
-  if (getModuleDebugInfo())
+  if (DebugInfo)
     // 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).
@@ -653,14 +549,6 @@ void CodeGenModule::Release() {
     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",
@@ -681,49 +569,14 @@ 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.FP32DenormalMode.Output !=
-                                  llvm::DenormalMode::IEEE);
+                              CodeGenOpts.FlushDenorm ? 1 : 0);
   }
 
-  if (LangOpts.EHAsynch)
-    getModule().addModuleFlag(llvm::Module::Warning, "eh-asynch", 1);
-
-  // Indicate whether this Module was compiled with -fopenmp
-  if (getLangOpts().OpenMP && !getLangOpts().OpenMPSimd)
-    getModule().addModuleFlag(llvm::Module::Max, "openmp", LangOpts.OpenMP);
-  if (getLangOpts().OpenMPIsDevice)
-    getModule().addModuleFlag(llvm::Module::Max, "openmp-device",
-                              LangOpts.OpenMP);
-
   // Emit OpenCL specific module metadata: OpenCL/SPIR version.
   if (LangOpts.OpenCL) {
     EmitOpenCLMetadata();
@@ -768,20 +621,6 @@ 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();
 
@@ -791,8 +630,8 @@ void CodeGenModule::Release() {
   if (getCodeGenOpts().EmitGcovArcs || getCodeGenOpts().EmitGcovNotes)
     EmitCoverageFile();
 
-  if (CGDebugInfo *DI = getModuleDebugInfo())
-    DI->finalize();
+  if (DebugInfo)
+    DebugInfo->finalize();
 
   if (getCodeGenOpts().EmitVersionIdentMetadata)
     EmitVersionIdentMetadata();
@@ -800,26 +639,7 @@ void CodeGenModule::Release() {
   if (!getCodeGenOpts().RecordCommandLine.empty())
     EmitCommandLineMetadata();
 
-  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);
-  if (getCodeGenOpts().StackAlignment)
-    getModule().setOverrideStackAlignment(getCodeGenOpts().StackAlignment);
-
-  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());
+  EmitTargetMetadata();
 }
 
 void CodeGenModule::EmitOpenCLMetadata() {
@@ -839,19 +659,6 @@ 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);
@@ -871,19 +678,6 @@ 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);
 }
 
@@ -1011,13 +805,8 @@ static bool shouldAssumeDSOLocal(const CodeGenModule &CGM,
     // In MinGW, variables without DLLImport can still be automatically
     // imported from a DLL by the linker; don't mark variables that
     // potentially could come from another DLL as DSO local.
-
-    // With EmulatedTLS, TLS variables can be autoimported from other DLLs
-    // (and this actually happens in the public interface of libstdc++), so
-    // such variables can't be marked as DSO local. (Native TLS variables
-    // can't be dllimported at all, though.)
     if (GV->isDeclarationForLinker() && isa<llvm::GlobalVariable>(GV) &&
-        (!GV->isThreadLocal() || CGM.getCodeGenOpts().EmulatedTLS))
+        !GV->isThreadLocal())
       return false;
   }
 
@@ -1043,17 +832,8 @@ 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) {
-    // On ELF, if -fno-semantic-interposition is specified and the target
-    // supports local aliases, there will be neither CC1
-    // -fsemantic-interposition nor -fhalf-no-semantic-interposition. Set
-    // dso_local on the function if using a local alias is preferable (can avoid
-    // PLT indirection).
-    if (!(isa<llvm::Function>(GV) && GV->canBenefitFromLocalAlias()))
-      return false;
-    return !(CGM.getLangOpts().SemanticInterposition ||
-             CGM.getLangOpts().HalfNoSemanticInterposition);
-  }
+  if (RM != llvm::Reloc::Static && !LOpts.PIE)
+    return false;
 
   // A definition cannot be preempted from an executable.
   if (!GV->isDeclarationForLinker())
@@ -1065,33 +845,25 @@ static bool shouldAssumeDSOLocal(const CodeGenModule &CGM,
   if (RM == llvm::Reloc::PIC_ && GV->hasExternalWeakLinkage())
     return false;
 
-  // PowerPC64 prefers TOC indirection to avoid copy relocations.
-  if (TT.isPPC64())
+  // 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)
     return false;
 
-  if (CGOpts.DirectAccessExternalData) {
-    // If -fdirect-access-external-data (default for -fno-pic), set dso_local
-    // for non-thread-local variables. If the symbol is not defined in the
-    // executable, a copy relocation will be needed at link time. dso_local is
-    // excluded for thread-local variables because they generally don't support
-    // copy relocations.
-    if (auto *Var = dyn_cast<llvm::GlobalVariable>(GV))
-      if (!Var->isThreadLocal())
-        return true;
-
-    // -fno-pic sets dso_local on a function declaration to allow direct
-    // accesses when taking its address (similar to a data symbol). If the
-    // function is not defined in the executable, a canonical PLT entry will be
-    // needed at link time. -fno-direct-access-external-data can avoid the
-    // canonical PLT entry. We don't generalize this condition to -fpie/-fpic as
-    // it could just cause trouble without providing perceptible benefits.
-    if (isa<llvm::Function>(GV) && !CGOpts.NoPLT && RM == llvm::Reloc::Static)
+  // If we can use copy relocations we can assume it is local.
+  if (auto *Var = dyn_cast<llvm::GlobalVariable>(GV))
+    if (!Var->isThreadLocal() &&
+        (RM == llvm::Reloc::Static || CGOpts.PIECopyRelocations))
       return true;
-  }
 
-  // If we can use copy relocations we can assume it is local.
+  // If we can use a plt entry as the symbol address we can assume it
+  // is local.
+  // FIXME: This should work for PIE, but the gold linker doesn't support it.
+  if (isa<llvm::Function>(GV) && !CGOpts.NoPLT && RM == llvm::Reloc::Static)
+    return true;
 
-  // Otherwise don't assume it is local.
+  // Otherwise don't assue it is local.
   return false;
 }
 
@@ -1147,9 +919,9 @@ static llvm::GlobalVariable::ThreadLocalMode GetLLVMTLSModel(StringRef S) {
       .Case("local-exec", llvm::GlobalVariable::LocalExecTLSModel);
 }
 
-llvm::GlobalVariable::ThreadLocalMode
-CodeGenModule::GetDefaultLLVMTLSModel() const {
-  switch (CodeGenOpts.getDefaultTLSModel()) {
+static llvm::GlobalVariable::ThreadLocalMode GetLLVMTLSModel(
+    CodeGenOptions::TLSModel M) {
+  switch (M) {
   case CodeGenOptions::GeneralDynamicTLSModel:
     return llvm::GlobalVariable::GeneralDynamicTLSModel;
   case CodeGenOptions::LocalDynamicTLSModel:
@@ -1166,7 +938,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 = GetDefaultLLVMTLSModel();
+  TLM = GetLLVMTLSModel(CodeGenOpts.getDefaultTLSModel());
 
   // Override the TLS model if it is explicitly specified.
   if (const TLSModelAttr *Attr = D.getAttr<TLSModelAttr>()) {
@@ -1226,56 +998,34 @@ static void AppendTargetMangling(const CodeGenModule &CGM,
   }
 }
 
-// Returns true if GD is a function decl with internal linkage and
-// needs a unique suffix after the mangled name.
-static bool isUniqueInternalLinkageDecl(GlobalDecl GD,
-                                        CodeGenModule &CGM) {
-  const Decl *D = GD.getDecl();
-  return !CGM.getModuleNameHash().empty() && isa<FunctionDecl>(D) &&
-         (CGM.getFunctionLinkage(GD) == llvm::GlobalValue::InternalLinkage);
-}
-
-static std::string getMangledNameImpl(CodeGenModule &CGM, GlobalDecl GD,
+static std::string getMangledNameImpl(const 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 (!CGM.getModuleNameHash().empty())
-    MC.needsUniqueInternalLinkageNames();
-  bool ShouldMangle = MC.shouldMangleDeclName(ND);
-  if (ShouldMangle)
-    MC.mangleName(GD.getWithDecl(ND), Out);
-  else {
+  if (MC.shouldMangleDeclName(ND)) {
+    llvm::raw_svector_ostream Out(Buffer);
+    if (const auto *D = dyn_cast<CXXConstructorDecl>(ND))
+      MC.mangleCXXCtor(D, GD.getCtorType(), Out);
+    else if (const auto *D = dyn_cast<CXXDestructorDecl>(ND))
+      MC.mangleCXXDtor(D, GD.getDtorType(), Out);
+    else
+      MC.mangleName(ND, Out);
+  } else {
     IdentifierInfo *II = ND->getIdentifier();
     assert(II && "Attempt to mangle unnamed decl.");
     const auto *FD = dyn_cast<FunctionDecl>(ND);
 
     if (FD &&
         FD->getType()->castAs<FunctionType>()->getCallConv() == CC_X86RegCall) {
+      llvm::raw_svector_ostream Out(Buffer);
       Out << "__regcall3__" << II->getName();
-    } else if (FD && FD->hasAttr<CUDAGlobalAttr>() &&
-               GD.getKernelReferenceKind() == KernelReferenceKind::Stub) {
-      Out << "__device_stub__" << II->getName();
     } else {
       Out << II->getName();
     }
   }
 
-  // Check if the module name hash should be appended for internal linkage
-  // symbols.   This should come before multi-version target suffixes are
-  // appended. This is to keep the name and module hash suffix of the
-  // internal linkage function together.  The unique suffix should only be
-  // added when name mangling is done to make sure that the final name can
-  // be properly demangled.  For example, for C functions without prototypes,
-  // name mangling is not done and the unique suffix should not be appeneded
-  // then.
-  if (ShouldMangle && isUniqueInternalLinkageDecl(GD, CGM)) {
-    assert(CGM.getCodeGenOpts().UniqueInternalLinkageNames &&
-           "Hash computed when not explicitly requested");
-    Out << CGM.getModuleNameHash();
-  }
-
   if (const auto *FD = dyn_cast<FunctionDecl>(ND))
     if (FD->isMultiVersion() && !OmitMultiVersionMangling) {
       switch (FD->getMultiVersionKind()) {
@@ -1293,12 +1043,7 @@ static std::string getMangledNameImpl(CodeGenModule &CGM, GlobalDecl GD,
       }
     }
 
-  // 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());
+  return Out.str();
 }
 
 void CodeGenModule::UpdateMultiVersionNames(GlobalDecl GD,
@@ -1355,40 +1100,19 @@ StringRef CodeGenModule::getMangledName(GlobalDecl GD) {
     }
   }
 
-  // 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;
-  }
+  auto FoundName = MangledDeclNames.find(CanonicalGD);
+  if (FoundName != MangledDeclNames.end())
+    return FoundName->second;
 
   // Keep the first result in the case of a mangling collision.
   const auto *ND = cast<NamedDecl>(GD.getDecl());
   std::string MangledName = getMangledNameImpl(*this, GD, ND);
 
-  // Ensure either we have different ABIs between host and device compilations,
-  // says host compilation following MSVC ABI but device compilation follows
-  // Itanium C++ ABI or, if they follow the same ABI, kernel names after
-  // mangling should be the same after name stubbing. The later checking is
-  // very important as the device kernel name being mangled in host-compilation
-  // is used to resolve the device binaries to be executed. Inconsistent naming
-  // result in undefined behavior. Even though we cannot check that naming
-  // directly between host- and device-compilations, the host- and
-  // device-mangling in host compilation could help catching certain ones.
-  assert(!isa<FunctionDecl>(ND) || !ND->hasAttr<CUDAGlobalAttr>() ||
-         getLangOpts().CUDAIsDevice ||
-         (getContext().getAuxTargetInfo() &&
-          (getContext().getAuxTargetInfo()->getCXXABI() !=
-           getContext().getTargetInfo().getCXXABI())) ||
-         getCUDARuntime().getDeviceSideName(ND) ==
-             getMangledNameImpl(
-                 *this,
-                 GD.getWithKernelReferenceKind(KernelReferenceKind::Kernel),
-                 ND));
+  // Adjust kernel stub mangling as we may need to be able to differentiate
+  // them from the kernel itself (e.g., for HIP).
+  if (auto *FD = dyn_cast<FunctionDecl>(GD.getDecl()))
+    if (!getLangOpts().CUDAIsDevice && FD->hasAttr<CUDAGlobalAttr>())
+      MangledName = getCUDARuntime().getDeviceStubName(MangledName);
 
   auto Result = Manglings.insert(std::make_pair(MangledName, GD));
   return MangledDeclNames[CanonicalGD] = Result.first->first();
@@ -1429,10 +1153,8 @@ 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,
-                                  bool IsDtorAttrFunc) {
-  if (CodeGenOpts.RegisterGlobalDtorsWithAtExit &&
-      (!getContext().getTargetInfo().getTriple().isOSAIX() || IsDtorAttrFunc)) {
+void CodeGenModule::AddGlobalDtor(llvm::Function *Dtor, int Priority) {
+  if (CodeGenOpts.RegisterGlobalDtorsWithAtExit) {
     DtorsUsingAtExit[Priority].push_back(Dtor);
     return;
   }
@@ -1509,11 +1231,10 @@ llvm::ConstantInt *CodeGenModule::CreateCrossDsoCfiTypeId(llvm::Metadata *MD) {
 
 void CodeGenModule::SetLLVMFunctionAttributes(GlobalDecl GD,
                                               const CGFunctionInfo &Info,
-                                              llvm::Function *F, bool IsThunk) {
+                                              llvm::Function *F) {
   unsigned CallingConv;
   llvm::AttributeList PAL;
-  ConstructAttributeList(F->getName(), Info, GD, PAL, CallingConv,
-                         /*AttrOnCallSite=*/false, IsThunk);
+  ConstructAttributeList(F->getName(), Info, GD, PAL, CallingConv, false);
   F->setAttributes(PAL);
   F->setCallingConv(static_cast<llvm::CallingConv::ID>(CallingConv));
 }
@@ -1546,18 +1267,10 @@ 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_device:
-    return 5;
-  case LangAS::opencl_global_host:
-    return 6;
+  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.
   default:
     return 0; // Assume private.
   }
@@ -1598,39 +1311,6 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
       QualType ty = parm->getType();
       std::string typeQuals;
 
-      // Get image and pipe access qualifier:
-      if (ty->isImageType() || ty->isPipeType()) {
-        const Decl *PDecl = parm;
-        if (auto *TD = dyn_cast<TypedefType>(ty))
-          PDecl = TD->getDecl();
-        const OpenCLAccessAttr *A = PDecl->getAttr<OpenCLAccessAttr>();
-        if (A && A->isWriteOnly())
-          accessQuals.push_back(llvm::MDString::get(VMContext, "write_only"));
-        else if (A && A->isReadWrite())
-          accessQuals.push_back(llvm::MDString::get(VMContext, "read_write"));
-        else
-          accessQuals.push_back(llvm::MDString::get(VMContext, "read_only"));
-      } else
-        accessQuals.push_back(llvm::MDString::get(VMContext, "none"));
-
-      // Get argument name.
-      argNames.push_back(llvm::MDString::get(VMContext, parm->getName()));
-
-      auto getTypeSpelling = [&](QualType Ty) {
-        auto typeName = Ty.getUnqualifiedType().getAsString(Policy);
-
-        if (Ty.isCanonical()) {
-          StringRef typeNameRef = typeName;
-          // Turn "unsigned type" to "utype"
-          if (typeNameRef.consume_front("unsigned "))
-            return std::string("u") + typeNameRef.str();
-          if (typeNameRef.consume_front("signed "))
-            return typeNameRef.str();
-        }
-
-        return typeName;
-      };
-
       if (ty->isPointerType()) {
         QualType pointeeTy = ty->getPointeeType();
 
@@ -1640,10 +1320,26 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
                 ArgInfoAddressSpace(pointeeTy.getAddressSpace()))));
 
         // Get argument type name.
-        std::string typeName = getTypeSpelling(pointeeTy) + "*";
-        std::string baseTypeName =
-            getTypeSpelling(pointeeTy.getCanonicalType()) + "*";
+        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 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);
+
         argBaseTypeNames.push_back(
             llvm::MDString::get(VMContext, baseTypeName));
 
@@ -1665,9 +1361,30 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
             llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(AddrSpc)));
 
         // Get argument type name.
-        ty = isPipe ? ty->castAs<PipeType>()->getElementType() : ty;
-        std::string typeName = getTypeSpelling(ty);
-        std::string baseTypeName = getTypeSpelling(ty.getCanonicalType());
+        std::string typeName;
+        if (isPipe)
+          typeName = ty.getCanonicalType()
+                         ->getAs<PipeType>()
+                         ->getElementType()
+                         .getAsString(Policy);
+        else
+          typeName = ty.getUnqualifiedType().getAsString(Policy);
+
+        // Turn "unsigned type" to "utype"
+        std::string::size_type pos = typeName.find("unsigned");
+        if (ty.isCanonical() && pos != std::string::npos)
+          typeName.erase(pos + 1, 8);
+
+        std::string baseTypeName;
+        if (isPipe)
+          baseTypeName = ty.getCanonicalType()
+                             ->getAs<PipeType>()
+                             ->getElementType()
+                             .getCanonicalType()
+                             .getAsString(Policy);
+        else
+          baseTypeName =
+              ty.getUnqualifiedType().getCanonicalType().getAsString(Policy);
 
         // Remove access qualifiers on images
         // (as they are inseparable from type in clang implementation,
@@ -1679,13 +1396,38 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
         }
 
         argTypeNames.push_back(llvm::MDString::get(VMContext, typeName));
+
+        // Turn "unsigned type" to "utype"
+        pos = baseTypeName.find("unsigned");
+        if (pos != std::string::npos)
+          baseTypeName.erase(pos + 1, 8);
+
         argBaseTypeNames.push_back(
             llvm::MDString::get(VMContext, baseTypeName));
 
         if (isPipe)
           typeQuals = "pipe";
       }
+
       argTypeQuals.push_back(llvm::MDString::get(VMContext, typeQuals));
+
+      // Get image and pipe access qualifier:
+      if (ty->isImageType() || ty->isPipeType()) {
+        const Decl *PDecl = parm;
+        if (auto *TD = dyn_cast<TypedefType>(ty))
+          PDecl = TD->getDecl();
+        const OpenCLAccessAttr *A = PDecl->getAttr<OpenCLAccessAttr>();
+        if (A && A->isWriteOnly())
+          accessQuals.push_back(llvm::MDString::get(VMContext, "write_only"));
+        else if (A && A->isReadWrite())
+          accessQuals.push_back(llvm::MDString::get(VMContext, "read_write"));
+        else
+          accessQuals.push_back(llvm::MDString::get(VMContext, "read_only"));
+      } else
+        accessQuals.push_back(llvm::MDString::get(VMContext, "none"));
+
+      // Get argument name.
+      argNames.push_back(llvm::MDString::get(VMContext, parm->getName()));
     }
 
   Fn->setMetadata("kernel_arg_addr_space",
@@ -1758,9 +1500,6 @@ 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);
 
@@ -1862,8 +1601,7 @@ void CodeGenModule::SetLLVMFunctionAttributesForDefinition(const Decl *D,
         B.addAttribute(llvm::Attribute::OptimizeForSize);
       B.addAttribute(llvm::Attribute::Cold);
     }
-    if (D->hasAttr<HotAttr>())
-      B.addAttribute(llvm::Attribute::Hot);
+
     if (D->hasAttr<MinSizeAttr>())
       B.addAttribute(llvm::Attribute::MinSize);
   }
@@ -1913,15 +1651,6 @@ void CodeGenModule::SetLLVMFunctionAttributesForDefinition(const Decl *D,
   }
 }
 
-void CodeGenModule::setLLVMFunctionFEnvAttributes(const FunctionDecl *D,
-                                                  llvm::Function *F) {
-  if (D->hasAttr<StrictFPAttr>()) {
-    llvm::AttrBuilder FuncAttrs;
-    FuncAttrs.addAttribute("strictfp");
-    F->addAttributes(llvm::AttributeList::FunctionIndex, FuncAttrs);
-  }
-}
-
 void CodeGenModule::SetCommonAttributes(GlobalDecl GD, llvm::GlobalValue *GV) {
   const Decl *D = GD.getDecl();
   if (dyn_cast_or_null<NamedDecl>(D))
@@ -1930,13 +1659,13 @@ void CodeGenModule::SetCommonAttributes(GlobalDecl GD, llvm::GlobalValue *GV) {
     GV->setVisibility(llvm::GlobalValue::DefaultVisibility);
 
   if (D && D->hasAttr<UsedAttr>())
-    addUsedOrCompilerUsedGlobal(GV);
+    addUsedGlobal(GV);
 
   if (CodeGenOpts.KeepStaticConsts && D && isa<VarDecl>(D)) {
     const auto *VD = cast<VarDecl>(D);
     if (VD->getType().isConstQualified() &&
         VD->getStorageDuration() == SD_Static)
-      addUsedOrCompilerUsedGlobal(GV);
+      addUsedGlobal(GV);
   }
 }
 
@@ -1946,7 +1675,6 @@ bool CodeGenModule::GetCPUAndFeaturesAttributes(GlobalDecl GD,
   // we have a decl for the function and it has a target attribute then
   // parse that and add it to the feature set.
   StringRef TargetCPU = getTarget().getTargetOpts().CPU;
-  StringRef TuneCPU = getTarget().getTargetOpts().TuneCPU;
   std::vector<std::string> Features;
   const auto *FD = dyn_cast_or_null<FunctionDecl>(GD.getDecl());
   FD = FD ? FD->getMostRecentDecl() : FD;
@@ -1967,14 +1695,9 @@ bool CodeGenModule::GetCPUAndFeaturesAttributes(GlobalDecl GD,
     // the function.
     if (TD) {
       ParsedTargetAttr ParsedAttr = TD->parse();
-      if (!ParsedAttr.Architecture.empty() &&
-          getTarget().isValidCPUName(ParsedAttr.Architecture)) {
+      if (ParsedAttr.Architecture != "" &&
+          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
@@ -1982,14 +1705,10 @@ bool CodeGenModule::GetCPUAndFeaturesAttributes(GlobalDecl GD,
     Features = getTarget().getTargetOpts().Features;
   }
 
-  if (!TargetCPU.empty()) {
+  if (TargetCPU != "") {
     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, ","));
@@ -2006,8 +1725,6 @@ void CodeGenModule::setNonAliasAttributes(GlobalDecl GD,
 
   if (D) {
     if (auto *GV = dyn_cast<llvm::GlobalVariable>(GO)) {
-      if (D->hasAttr<RetainAttr>())
-        addUsedGlobal(GV);
       if (auto *SA = D->getAttr<PragmaClangBSSSectionAttr>())
         GV->addAttribute("bss-section", SA->getName());
       if (auto *SA = D->getAttr<PragmaClangDataSectionAttr>())
@@ -2019,8 +1736,6 @@ void CodeGenModule::setNonAliasAttributes(GlobalDecl GD,
     }
 
     if (auto *F = dyn_cast<llvm::Function>(GO)) {
-      if (D->hasAttr<RetainAttr>())
-        addUsedGlobal(F);
       if (auto *SA = D->getAttr<PragmaClangTextSectionAttr>())
         if (!D->getAttr<SectionAttr>())
           F->addFnAttr("implicit-section-name", SA->getName());
@@ -2030,11 +1745,8 @@ 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.
-        llvm::AttrBuilder RemoveAttrs;
-        RemoveAttrs.addAttribute("target-cpu");
-        RemoveAttrs.addAttribute("target-features");
-        RemoveAttrs.addAttribute("tune-cpu");
-        F->removeAttributes(llvm::AttributeList::FunctionIndex, RemoveAttrs);
+        F->removeFnAttr("target-cpu");
+        F->removeFnAttr("target-features");
         F->addAttributes(llvm::AttributeList::FunctionIndex, Attrs);
       }
     }
@@ -2052,7 +1764,7 @@ void CodeGenModule::SetInternalFunctionAttributes(GlobalDecl GD,
                                                   llvm::Function *F,
                                                   const CGFunctionInfo &FI) {
   const Decl *D = GD.getDecl();
-  SetLLVMFunctionAttributes(GD, FI, F, /*IsThunk=*/false);
+  SetLLVMFunctionAttributes(GD, FI, F);
   SetLLVMFunctionAttributesForDefinition(D, F);
 
   F->setLinkage(llvm::Function::InternalLinkage);
@@ -2106,8 +1818,7 @@ void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F,
   const auto *FD = cast<FunctionDecl>(GD.getDecl());
 
   if (!IsIncompleteFunction)
-    SetLLVMFunctionAttributes(GD, getTypes().arrangeGlobalDeclaration(GD), F,
-                              IsThunk);
+    SetLLVMFunctionAttributes(GD, getTypes().arrangeGlobalDeclaration(GD), F);
 
   // Add the Returned attribute for "this", except for iOS 5 and earlier
   // where substantial code, including the libstdc++ dylib, was compiled with
@@ -2136,16 +1847,9 @@ void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F,
   else if (const auto *SA = FD->getAttr<SectionAttr>())
      F->setSection(SA->getName());
 
-  // If we plan on emitting this inline builtin, we can't treat it as a builtin.
   if (FD->isInlineBuiltinDeclaration()) {
-    const FunctionDecl *FDBody;
-    bool HasBody = FD->hasBody(FDBody);
-    (void)HasBody;
-    assert(HasBody && "Inline builtin declarations should always have an "
-                      "available body!");
-    if (shouldEmitFunction(FDBody))
-      F->addAttribute(llvm::AttributeList::FunctionIndex,
-                      llvm::Attribute::NoBuiltin);
+    F->addAttribute(llvm::AttributeList::FunctionIndex,
+                    llvm::Attribute::NoBuiltin);
   }
 
   if (FD->isReplaceableGlobalAllocationFunction()) {
@@ -2153,6 +1857,15 @@ void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F,
     // default, only if it is invoked by a new-expression or delete-expression.
     F->addAttribute(llvm::AttributeList::FunctionIndex,
                     llvm::Attribute::NoBuiltin);
+
+    // A sane operator new returns a non-aliasing pointer.
+    // FIXME: Also add NonNull attribute to the return value
+    // for the non-nothrow forms?
+    auto Kind = FD->getDeclName().getCXXOverloadedOperator();
+    if (getCodeGenOpts().AssumeSaneOperatorNew &&
+        (Kind == OO_New || Kind == OO_Array_New))
+      F->addAttribute(llvm::AttributeList::ReturnIndex,
+                      llvm::Attribute::NoAlias);
   }
 
   if (isa<CXXConstructorDecl>(FD) || isa<CXXDestructorDecl>(FD))
@@ -2191,7 +1904,7 @@ void CodeGenModule::SetFunctionAttributes(GlobalDecl GD, llvm::Function *F,
 }
 
 void CodeGenModule::addUsedGlobal(llvm::GlobalValue *GV) {
-  assert((isa<llvm::Function>(GV) || !GV->isDeclaration()) &&
+  assert(!GV->isDeclaration() &&
          "Only globals with definition can force usage.");
   LLVMUsed.emplace_back(GV);
 }
@@ -2202,15 +1915,6 @@ void CodeGenModule::addCompilerUsedGlobal(llvm::GlobalValue *GV) {
   LLVMCompilerUsed.emplace_back(GV);
 }
 
-void CodeGenModule::addUsedOrCompilerUsedGlobal(llvm::GlobalValue *GV) {
-  assert((isa<llvm::Function>(GV) || !GV->isDeclaration()) &&
-         "Only globals with definition can force usage.");
-  if (getTriple().isOSBinFormatELF())
-    LLVMCompilerUsed.emplace_back(GV);
-  else
-    LLVMUsed.emplace_back(GV);
-}
-
 static void emitUsed(CodeGenModule &CGM, StringRef Name,
                      std::vector<llvm::WeakTrackingVH> &List) {
   // Don't create llvm.used if there is no need.
@@ -2406,13 +2110,6 @@ 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;
@@ -2528,47 +2225,13 @@ llvm::Constant *CodeGenModule::EmitAnnotationLineNo(SourceLocation L) {
   return llvm::ConstantInt::get(Int32Ty, LineNo);
 }
 
-llvm::Constant *CodeGenModule::EmitAnnotationArgs(const AnnotateAttr *Attr) {
-  ArrayRef<Expr *> Exprs = {Attr->args_begin(), Attr->args_size()};
-  if (Exprs.empty())
-    return llvm::ConstantPointerNull::get(Int8PtrTy);
-
-  llvm::FoldingSetNodeID ID;
-  for (Expr *E : Exprs) {
-    ID.Add(cast<clang::ConstantExpr>(E)->getAPValueResult());
-  }
-  llvm::Constant *&Lookup = AnnotationArgs[ID.ComputeHash()];
-  if (Lookup)
-    return Lookup;
-
-  llvm::SmallVector<llvm::Constant *, 4> LLVMArgs;
-  LLVMArgs.reserve(Exprs.size());
-  ConstantEmitter ConstEmiter(*this);
-  llvm::transform(Exprs, std::back_inserter(LLVMArgs), [&](const Expr *E) {
-    const auto *CE = cast<clang::ConstantExpr>(E);
-    return ConstEmiter.emitAbstract(CE->getBeginLoc(), CE->getAPValueResult(),
-                                    CE->getType());
-  });
-  auto *Struct = llvm::ConstantStruct::getAnon(LLVMArgs);
-  auto *GV = new llvm::GlobalVariable(getModule(), Struct->getType(), true,
-                                      llvm::GlobalValue::PrivateLinkage, Struct,
-                                      ".args");
-  GV->setSection(AnnotationSection);
-  GV->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
-  auto *Bitcasted = llvm::ConstantExpr::getBitCast(GV, Int8PtrTy);
-
-  Lookup = Bitcasted;
-  return Bitcasted;
-}
-
 llvm::Constant *CodeGenModule::EmitAnnotateAttr(llvm::GlobalValue *GV,
                                                 const AnnotateAttr *AA,
                                                 SourceLocation L) {
   // Get the globals for file name, annotation, and the line number.
   llvm::Constant *AnnoGV = EmitAnnotationString(AA->getAnnotation()),
                  *UnitGV = EmitAnnotationUnit(L),
-                 *LineNoCst = EmitAnnotationLineNo(L),
-                 *Args = EmitAnnotationArgs(AA);
+                 *LineNoCst = EmitAnnotationLineNo(L);
 
   llvm::Constant *ASZeroGV = GV;
   if (GV->getAddressSpace() != 0) {
@@ -2577,12 +2240,11 @@ llvm::Constant *CodeGenModule::EmitAnnotateAttr(llvm::GlobalValue *GV,
   }
 
   // Create the ConstantStruct for the global annotation.
-  llvm::Constant *Fields[] = {
-      llvm::ConstantExpr::getBitCast(ASZeroGV, Int8PtrTy),
-      llvm::ConstantExpr::getBitCast(AnnoGV, Int8PtrTy),
-      llvm::ConstantExpr::getBitCast(UnitGV, Int8PtrTy),
-      LineNoCst,
-      Args,
+  llvm::Constant *Fields[4] = {
+    llvm::ConstantExpr::getBitCast(ASZeroGV, Int8PtrTy),
+    llvm::ConstantExpr::getBitCast(AnnoGV, Int8PtrTy),
+    llvm::ConstantExpr::getBitCast(UnitGV, Int8PtrTy),
+    LineNoCst
   };
   return llvm::ConstantStruct::getAnon(Fields);
 }
@@ -2595,28 +2257,29 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D,
     Annotations.push_back(EmitAnnotateAttr(GV, I, D->getLocation()));
 }
 
-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()))
+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()))
     return true;
-  // NoSanitize by location.
+  // Blacklist by location.
   if (Loc.isValid())
-    return NoSanitizeL.containsLocation(Kind, Loc);
+    return SanitizerBL.isBlacklistedLocation(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 NoSanitizeL.containsFile(Kind, MainFile->getName());
+    return SanitizerBL.isBlacklistedFile(Kind, MainFile->getName());
   }
   return false;
 }
 
-bool CodeGenModule::isInNoSanitizeList(llvm::GlobalVariable *GV,
-                                       SourceLocation Loc, QualType Ty,
-                                       StringRef Category) const {
-  // For now globals can be ignored only in ASan and KASan.
+bool CodeGenModule::isInSanitizerBlacklist(llvm::GlobalVariable *GV,
+                                           SourceLocation Loc, QualType Ty,
+                                           StringRef Category) const {
+  // For now globals can be blacklisted only in ASan and KASan.
   const SanitizerMask EnabledAsanMask =
       LangOpts.Sanitize.Mask &
       (SanitizerKind::Address | SanitizerKind::KernelAddress |
@@ -2624,22 +2287,22 @@ bool CodeGenModule::isInNoSanitizeList(llvm::GlobalVariable *GV,
        SanitizerKind::MemTag);
   if (!EnabledAsanMask)
     return false;
-  const auto &NoSanitizeL = getContext().getNoSanitizeList();
-  if (NoSanitizeL.containsGlobal(EnabledAsanMask, GV->getName(), Category))
+  const auto &SanitizerBL = getContext().getSanitizerBlacklist();
+  if (SanitizerBL.isBlacklistedGlobal(EnabledAsanMask, GV->getName(), Category))
     return true;
-  if (NoSanitizeL.containsLocation(EnabledAsanMask, Loc, Category))
+  if (SanitizerBL.isBlacklistedLocation(EnabledAsanMask, Loc, Category))
     return true;
   // Check global type.
   if (!Ty.isNull()) {
     // Drill down the array types: if global variable of a fixed type is
-    // not sanitized, we also don't instrument arrays of them.
+    // blacklisted, we also don't instrument arrays of them.
     while (auto AT = dyn_cast<ArrayType>(Ty.getTypePtr()))
       Ty = AT->getElementType();
     Ty = Ty.getCanonicalType().getUnqualifiedType();
-    // Only record types (classes, structs etc.) are ignored.
+    // We allow to blacklist only record types (classes, structs etc.)
     if (Ty->isRecordType()) {
       std::string TypeStr = Ty.getAsString(getContext().getPrintingPolicy());
-      if (NoSanitizeL.containsType(EnabledAsanMask, TypeStr, Category))
+      if (SanitizerBL.isBlacklistedType(EnabledAsanMask, TypeStr, Category))
         return true;
     }
   }
@@ -2672,34 +2335,6 @@ bool CodeGenModule::imbueXRayAttrs(llvm::Function *Fn, SourceLocation Loc,
   return true;
 }
 
-bool CodeGenModule::isProfileInstrExcluded(llvm::Function *Fn,
-                                           SourceLocation Loc) const {
-  const auto &ProfileList = getContext().getProfileList();
-  // If the profile list is empty, then instrument everything.
-  if (ProfileList.isEmpty())
-    return false;
-  CodeGenOptions::ProfileInstrKind Kind = getCodeGenOpts().getProfileInstr();
-  // First, check the function name.
-  Optional<bool> V = ProfileList.isFunctionExcluded(Fn->getName(), Kind);
-  if (V.hasValue())
-    return *V;
-  // Next, check the source location.
-  if (Loc.isValid()) {
-    Optional<bool> V = ProfileList.isLocationExcluded(Loc, Kind);
-    if (V.hasValue())
-      return *V;
-  }
-  // If location is unknown, this may be a compiler-generated function. Assume
-  // it's located in the main file.
-  auto &SM = Context.getSourceManager();
-  if (const auto *MainFile = SM.getFileEntryForID(SM.getMainFileID())) {
-    Optional<bool> V = ProfileList.isFileExcluded(MainFile->getName(), Kind);
-    if (V.hasValue())
-      return *V;
-  }
-  return ProfileList.getDefault();
-}
-
 bool CodeGenModule::MustBeEmitted(const ValueDecl *Global) {
   // Never defer when EmitAllDecls is specified.
   if (LangOpts.EmitAllDecls)
@@ -2716,24 +2351,19 @@ bool CodeGenModule::MustBeEmitted(const ValueDecl *Global) {
 }
 
 bool CodeGenModule::MayBeEmittedEagerly(const ValueDecl *Global) {
-  // In OpenMP 5.0 variables and function may be marked as
-  // device_type(host/nohost) and we should not emit them eagerly unless we sure
-  // that they must be emitted on the host/device. To be sure we need to have
-  // seen a declare target with an explicit mentioning of the function, we know
-  // we have if the level of the declare target attribute is -1. Note that we
-  // check somewhere else if we should emit this at all.
-  if (LangOpts.OpenMP >= 50 && !LangOpts.OpenMPSimd) {
-    llvm::Optional<OMPDeclareTargetDeclAttr *> ActiveAttr =
-        OMPDeclareTargetDeclAttr::getActiveAttr(Global);
-    if (!ActiveAttr || (*ActiveAttr)->getLevel() != (unsigned)-1)
-      return false;
-  }
-
   if (const auto *FD = dyn_cast<FunctionDecl>(Global)) {
     if (FD->getTemplateSpecializationKind() == TSK_ImplicitInstantiation)
       // Implicit template instantiations may change linkage if they are later
       // explicitly instantiated, so they should not be emitted eagerly.
       return false;
+    // In OpenMP 5.0 function may be marked as device_type(nohost) and we should
+    // not emit them eagerly unless we sure that the function must be emitted on
+    // the host.
+    if (LangOpts.OpenMP >= 50 && !LangOpts.OpenMPSimd &&
+        !LangOpts.OpenMPIsDevice &&
+        !OMPDeclareTargetDeclAttr::getDeviceType(FD) &&
+        !FD->isUsed(/*CheckUsedAttr=*/false) && !FD->isReferenced())
+      return false;
   }
   if (const auto *VD = dyn_cast<VarDecl>(Global))
     if (Context.getInlineVariableDefinitionKind(VD) ==
@@ -2752,8 +2382,13 @@ bool CodeGenModule::MayBeEmittedEagerly(const ValueDecl *Global) {
   return true;
 }
 
-ConstantAddress CodeGenModule::GetAddrOfMSGuidDecl(const MSGuidDecl *GD) {
-  StringRef Name = getMangledName(GD);
+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(), '-', '_');
 
   // The UUID descriptor should be pointer aligned.
   CharUnits Alignment = CharUnits::fromQuantity(PointerAlignInBytes);
@@ -2762,30 +2397,8 @@ ConstantAddress CodeGenModule::GetAddrOfMSGuidDecl(const MSGuidDecl *GD) {
   if (llvm::GlobalVariable *GV = getModule().getNamedGlobal(Name))
     return ConstantAddress(GV, Alignment);
 
-  ConstantEmitter Emitter(*this);
-  llvm::Constant *Init;
-
-  APValue &V = GD->getAsAPValue();
-  if (!V.isAbsent()) {
-    // If possible, emit the APValue version of the initializer. In particular,
-    // this gets the type of the constant right.
-    Init = Emitter.emitForInitializer(
-        GD->getAsAPValue(), GD->getType().getAddressSpace(), GD->getType());
-  } else {
-    // As a fallback, directly construct the constant.
-    // FIXME: This may get padding wrong under esoteric struct layout rules.
-    // MSVC appears to create a complete type 'struct __s_GUID' that it
-    // presumably uses to represent these constants.
-    MSGuidDecl::Parts Parts = GD->getParts();
-    llvm::Constant *Fields[4] = {
-        llvm::ConstantInt::get(Int32Ty, Parts.Part1),
-        llvm::ConstantInt::get(Int16Ty, Parts.Part2),
-        llvm::ConstantInt::get(Int16Ty, Parts.Part3),
-        llvm::ConstantDataArray::getRaw(
-            StringRef(reinterpret_cast<char *>(Parts.Part4And5), 8), 8,
-            Int8Ty)};
-    Init = llvm::ConstantStruct::getAnon(Fields);
-  }
+  llvm::Constant *Init = EmitUuidofInitializer(Uuid);
+  assert(Init && "failed to initialize as constant");
 
   auto *GV = new llvm::GlobalVariable(
       getModule(), Init->getType(),
@@ -2793,42 +2406,6 @@ ConstantAddress CodeGenModule::GetAddrOfMSGuidDecl(const MSGuidDecl *GD) {
   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);
 }
 
@@ -2853,7 +2430,9 @@ ConstantAddress CodeGenModule::GetWeakRefReference(const ValueDecl *VD) {
                                       GlobalDecl(cast<FunctionDecl>(VD)),
                                       /*ForVTable=*/false);
   else
-    Aliasee = GetOrCreateLLVMGlobal(AA->getAliasee(), DeclTy, 0, nullptr);
+    Aliasee = GetOrCreateLLVMGlobal(AA->getAliasee(),
+                                    llvm::PointerType::getUnqual(DeclTy),
+                                    nullptr);
 
   auto *F = cast<llvm::GlobalValue>(Aliasee);
   F->setLinkage(llvm::Function::ExternalWeakLinkage);
@@ -2889,8 +2468,7 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
           !Global->hasAttr<CUDAGlobalAttr>() &&
           !Global->hasAttr<CUDAConstantAttr>() &&
           !Global->hasAttr<CUDASharedAttr>() &&
-          !Global->getType()->isCUDADeviceBuiltinSurfaceType() &&
-          !Global->getType()->isCUDADeviceBuiltinTextureType())
+          !(LangOpts.HIP && Global->hasAttr<HIPPinnedShadowAttr>()))
         return;
     } else {
       // We need to emit host-side 'shadows' for all global
@@ -2983,6 +2561,11 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
     return;
   }
 
+    // Check if this must be emitted as declare variant.
+  if (LangOpts.OpenMP && isa<FunctionDecl>(Global) && OpenMPRuntime &&
+      OpenMPRuntime->emitDeclareVariant(GD, /*IsForDefinition=*/false))
+    return;
+
   // If we're deferring emission of a C++ variable with an
   // initializer, remember the order in which it appeared in the file.
   if (getLangOpts().CPlusPlus && isa<VarDecl>(Global) &&
@@ -3143,7 +2726,7 @@ bool CodeGenModule::shouldEmitFunction(GlobalDecl GD) {
   if (CodeGenOpts.OptimizationLevel == 0 && !F->hasAttr<AlwaysInlineAttr>())
     return false;
 
-  if (F->hasAttr<DLLImportAttr>() && !F->hasAttr<AlwaysInlineAttr>()) {
+  if (F->hasAttr<DLLImportAttr>()) {
     // Check whether it would be safe to inline this dllimport function.
     DLLImportFunctionVisitor Visitor;
     Visitor.TraverseFunctionDecl(const_cast<FunctionDecl*>(F));
@@ -3165,8 +2748,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 through asm label/`__builtin_` trickery is
-  // clearly not equivalent to the real implementation.
+  // but a function that calls itself is clearly not equivalent to the real
+  // implementation.
   // This happens in glibc's btowc and in some configure checks.
   return !isTriviallyRecursive(F);
 }
@@ -3188,6 +2771,50 @@ void CodeGenModule::EmitMultiVersionFunctionDefinition(GlobalDecl GD,
     EmitGlobalFunctionDefinition(GD, GV);
 }
 
+void CodeGenModule::emitOpenMPDeviceFunctionRedefinition(
+    GlobalDecl OldGD, GlobalDecl NewGD, llvm::GlobalValue *GV) {
+  assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
+         OpenMPRuntime && "Expected OpenMP device mode.");
+  const auto *D = cast<FunctionDecl>(OldGD.getDecl());
+
+  // Compute the function info and LLVM type.
+  const CGFunctionInfo &FI = getTypes().arrangeGlobalDeclaration(OldGD);
+  llvm::FunctionType *Ty = getTypes().GetFunctionType(FI);
+
+  // Get or create the prototype for the function.
+  if (!GV || (GV->getType()->getElementType() != Ty)) {
+    GV = cast<llvm::GlobalValue>(GetOrCreateLLVMFunction(
+        getMangledName(OldGD), Ty, GlobalDecl(), /*ForVTable=*/false,
+        /*DontDefer=*/true, /*IsThunk=*/false, llvm::AttributeList(),
+        ForDefinition));
+    SetFunctionAttributes(OldGD, cast<llvm::Function>(GV),
+                          /*IsIncompleteFunction=*/false,
+                          /*IsThunk=*/false);
+  }
+  // We need to set linkage and visibility on the function before
+  // generating code for it because various parts of IR generation
+  // want to propagate this information down (e.g. to local static
+  // declarations).
+  auto *Fn = cast<llvm::Function>(GV);
+  setFunctionLinkage(OldGD, Fn);
+
+  // FIXME: this is redundant with part of
+  // setFunctionDefinitionAttributes
+  setGVProperties(Fn, OldGD);
+
+  MaybeHandleStaticInExternC(D, Fn);
+
+  maybeSetTrivialComdat(*D, *Fn);
+
+  CodeGenFunction(*this).GenerateCode(NewGD, Fn, FI);
+
+  setNonAliasAttributes(OldGD, Fn);
+  SetLLVMFunctionAttributesForDefinition(D, Fn);
+
+  if (D->hasAttr<AnnotateAttr>())
+    AddGlobalAnnotations(D, Fn);
+}
+
 void CodeGenModule::EmitGlobalDefinition(GlobalDecl GD, llvm::GlobalValue *GV) {
   const auto *D = cast<ValueDecl>(GD.getDecl());
 
@@ -3253,9 +2880,7 @@ TargetMVPriority(const TargetInfo &TI,
 }
 
 void CodeGenModule::emitMultiVersionFunctions() {
-  std::vector<GlobalDecl> MVFuncsToEmit;
-  MultiVersionFuncs.swap(MVFuncsToEmit);
-  for (GlobalDecl GD : MVFuncsToEmit) {
+  for (GlobalDecl GD : MultiVersionFuncs) {
     SmallVector<CodeGenFunction::MultiVersionResolverOption, 10> Options;
     const FunctionDecl *FD = cast<FunctionDecl>(GD.getDecl());
     getContext().forEachMultiversionedFunctionVersion(
@@ -3309,17 +2934,6 @@ 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) {
@@ -3394,7 +3008,7 @@ void CodeGenModule::emitCPUDispatchDefinition(GlobalDecl GD) {
     ++Index;
   }
 
-  llvm::stable_sort(
+  llvm::sort(
       Options, [](const CodeGenFunction::MultiVersionResolverOption &LHS,
                   const CodeGenFunction::MultiVersionResolverOption &RHS) {
         return CodeGenFunction::GetX86CpuSupportsMask(LHS.Conditions.Features) >
@@ -3515,9 +3129,14 @@ 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()) {
-      if (FD->hasAttr<TargetAttr>())
+      const auto *TA = FD->getAttr<TargetAttr>();
+      if (TA && TA->isDefaultVersion())
         UpdateMultiVersionNames(GD, FD);
       if (!IsForDefinition)
         return GetOrCreateMultiVersionResolver(GD, Ty, FD);
@@ -3557,7 +3176,7 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction(
     }
 
     if ((isa<llvm::Function>(Entry) || isa<llvm::GlobalAlias>(Entry)) &&
-        (Entry->getValueType() == Ty)) {
+        (Entry->getType()->getElementType() == Ty)) {
       return Entry;
     }
 
@@ -3606,7 +3225,7 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction(
     }
 
     llvm::Constant *BC = llvm::ConstantExpr::getBitCast(
-        F, Entry->getValueType()->getPointerTo());
+        F, Entry->getType()->getElementType()->getPointerTo());
     addGlobalValReplacement(Entry, BC);
   }
 
@@ -3665,7 +3284,7 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction(
 
   // Make sure the result is of the requested type.
   if (!IsIncompleteFunction) {
-    assert(F->getFunctionType() == Ty);
+    assert(F->getType()->getElementType() == Ty);
     return F;
   }
 
@@ -3681,8 +3300,6 @@ llvm::Constant *CodeGenModule::GetAddrOfFunction(GlobalDecl GD,
                                                  bool ForVTable,
                                                  bool DontDefer,
                                               ForDefinition_t IsForDefinition) {
-  assert(!cast<FunctionDecl>(GD.getDecl())->isConsteval() &&
-         "consteval function should never be emitted");
   // If there was no specific requested type, just convert it now.
   if (!Ty) {
     const auto *FD = cast<FunctionDecl>(GD.getDecl());
@@ -3700,19 +3317,9 @@ llvm::Constant *CodeGenModule::GetAddrOfFunction(GlobalDecl GD,
   }
 
   StringRef MangledName = getMangledName(GD);
-  auto *F = GetOrCreateLLVMFunction(MangledName, Ty, GD, ForVTable, DontDefer,
-                                    /*IsThunk=*/false, llvm::AttributeList(),
-                                    IsForDefinition);
-  // Returns kernel handle for HIP kernel stub function.
-  if (LangOpts.CUDA && !LangOpts.CUDAIsDevice &&
-      cast<FunctionDecl>(GD.getDecl())->hasAttr<CUDAGlobalAttr>()) {
-    auto *Handle = getCUDARuntime().getKernelHandle(
-        cast<llvm::Function>(F->stripPointerCasts()), GD);
-    if (IsForDefinition)
-      return F;
-    return llvm::ConstantExpr::getBitCast(Handle, Ty->getPointerTo());
-  }
-  return F;
+  return GetOrCreateLLVMFunction(MangledName, Ty, GD, ForVTable, DontDefer,
+                                 /*IsThunk=*/false, llvm::AttributeList(),
+                                 IsForDefinition);
 }
 
 static const FunctionDecl *
@@ -3721,8 +3328,8 @@ GetRuntimeFunctionDecl(ASTContext &C, StringRef Name) {
   DeclContext *DC = TranslationUnitDecl::castToDeclContext(TUDecl);
 
   IdentifierInfo &CII = C.Idents.get(Name);
-  for (const auto *Result : DC->lookup(&CII))
-    if (const auto *FD = dyn_cast<FunctionDecl>(Result))
+  for (const auto &Result : DC->lookup(&CII))
+    if (const auto FD = dyn_cast<FunctionDecl>(Result))
       return FD;
 
   if (!C.getLangOpts().CPlusPlus)
@@ -3736,15 +3343,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)) {
-      const NamespaceDecl *ND = dyn_cast<NamespaceDecl>(Result);
-      if (auto *LSD = dyn_cast<LinkageSpecDecl>(Result))
-        for (const auto *Result : LSD->lookup(&NS))
+    for (const auto &Result : DC->lookup(&NS)) {
+      NamespaceDecl *ND = dyn_cast<NamespaceDecl>(Result);
+      if (auto LSD = dyn_cast<LinkageSpecDecl>(Result))
+        for (const auto &Result : LSD->lookup(&NS))
           if ((ND = dyn_cast<NamespaceDecl>(Result)))
             break;
 
       if (ND)
-        for (const auto *Result : ND->lookup(&CXXII))
+        for (const auto &Result : ND->lookup(&CXXII))
           if (const auto *FD = dyn_cast<FunctionDecl>(Result))
             return FD;
     }
@@ -3815,9 +3422,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 and address
-/// space. 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.  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.
@@ -3826,8 +3433,9 @@ 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::Type *Ty,
-                                     unsigned AddrSpace, const VarDecl *D,
+CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName,
+                                     llvm::PointerType *Ty,
+                                     const VarDecl *D,
                                      ForDefinition_t IsForDefinition) {
   // Lookup the entry, lazily creating it if necessary.
   llvm::GlobalValue *Entry = GetGlobalValue(MangledName);
@@ -3844,7 +3452,7 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName, llvm::Type *Ty,
     if (LangOpts.OpenMP && !LangOpts.OpenMPSimd && D)
       getOpenMPRuntime().registerTargetGlobalVariable(D, Entry);
 
-    if (Entry->getValueType() == Ty && Entry->getAddressSpace() == AddrSpace)
+    if (Entry->getType() == Ty)
       return Entry;
 
     // If there are two attempts to define the same mangled name, issue an
@@ -3868,24 +3476,22 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName, llvm::Type *Ty,
     }
 
     // Make sure the result is of the correct type.
-    if (Entry->getType()->getAddressSpace() != AddrSpace) {
-      return llvm::ConstantExpr::getAddrSpaceCast(Entry,
-                                                  Ty->getPointerTo(AddrSpace));
-    }
+    if (Entry->getType()->getAddressSpace() != Ty->getAddressSpace())
+      return llvm::ConstantExpr::getAddrSpaceCast(Entry, Ty);
 
     // (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->getPointerTo(AddrSpace));
+      return llvm::ConstantExpr::getBitCast(Entry, Ty);
   }
 
-  auto DAddrSpace = GetGlobalVarAddressSpace(D);
-  auto TargetAddrSpace = getContext().getTargetAddressSpace(DAddrSpace);
+  auto AddrSpace = GetGlobalVarAddressSpace(D);
+  auto TargetAddrSpace = getContext().getTargetAddressSpace(AddrSpace);
 
   auto *GV = new llvm::GlobalVariable(
-      getModule(), Ty, false, llvm::GlobalValue::ExternalLinkage, nullptr,
-      MangledName, nullptr, llvm::GlobalVariable::NotThreadLocal,
-      TargetAddrSpace);
+      getModule(), Ty->getElementType(), 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.
@@ -3969,7 +3575,7 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName, llvm::Type *Ty,
           llvm::Constant *Init = emitter.tryEmitForInitializer(*InitDecl);
           if (Init) {
             auto *InitType = Init->getType();
-            if (GV->getValueType() != InitType) {
+            if (GV->getType()->getElementType() != 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
@@ -3996,51 +3602,42 @@ CodeGenModule::GetOrCreateLLVMGlobal(StringRef MangledName, llvm::Type *Ty,
     }
   }
 
-  if (GV->isDeclaration()) {
+  if (GV->isDeclaration())
     getTargetCodeGenInfo().setTargetAttributes(D, GV, *this);
-    // External HIP managed variables needed to be recorded for transformation
-    // in both device and host compilations.
-    if (getLangOpts().CUDA && D && D->hasAttr<HIPManagedAttr>() &&
-        D->hasExternalStorage())
-      getCUDARuntime().handleVarRegistration(D, *GV);
-  }
 
   LangAS ExpectedAS =
       D ? D->getType().getAddressSpace()
         : (LangOpts.OpenCL ? LangAS::opencl_global : LangAS::Default);
-  assert(getContext().getTargetAddressSpace(ExpectedAS) == AddrSpace);
-  if (DAddrSpace != ExpectedAS) {
-    return getTargetCodeGenInfo().performAddrSpaceCast(
-        *this, GV, DAddrSpace, ExpectedAS, Ty->getPointerTo(AddrSpace));
-  }
+  assert(getContext().getTargetAddressSpace(ExpectedAS) ==
+         Ty->getPointerAddressSpace());
+  if (AddrSpace != ExpectedAS)
+    return getTargetCodeGenInfo().performAddrSpaceCast(*this, GV, AddrSpace,
+                                                       ExpectedAS, Ty);
 
   return GV;
 }
 
 llvm::Constant *
-CodeGenModule::GetAddrOfGlobal(GlobalDecl GD, ForDefinition_t IsForDefinition) {
+CodeGenModule::GetAddrOfGlobal(GlobalDecl GD,
+                               ForDefinition_t IsForDefinition) {
   const Decl *D = GD.getDecl();
-
   if (isa<CXXConstructorDecl>(D) || isa<CXXDestructorDecl>(D))
     return getAddrOfCXXStructor(GD, /*FnInfo=*/nullptr, /*FnType=*/nullptr,
                                 /*DontDefer=*/false, IsForDefinition);
-
-  if (isa<CXXMethodDecl>(D)) {
-    auto FInfo =
-        &getTypes().arrangeCXXMethodDeclaration(cast<CXXMethodDecl>(D));
+  else if (isa<CXXMethodDecl>(D)) {
+    auto FInfo = &getTypes().arrangeCXXMethodDeclaration(
+        cast<CXXMethodDecl>(D));
     auto Ty = getTypes().GetFunctionType(*FInfo);
     return GetAddrOfFunction(GD, Ty, /*ForVTable=*/false, /*DontDefer=*/false,
                              IsForDefinition);
-  }
-
-  if (isa<FunctionDecl>(D)) {
+  } else if (isa<FunctionDecl>(D)) {
     const CGFunctionInfo &FI = getTypes().arrangeGlobalDeclaration(GD);
     llvm::FunctionType *Ty = getTypes().GetFunctionType(FI);
     return GetAddrOfFunction(GD, Ty, /*ForVTable=*/false, /*DontDefer=*/false,
                              IsForDefinition);
-  }
-
-  return GetAddrOfGlobalVar(cast<VarDecl>(D), /*Ty=*/nullptr, IsForDefinition);
+  } else
+    return GetAddrOfGlobalVar(cast<VarDecl>(D), /*Ty=*/nullptr,
+                              IsForDefinition);
 }
 
 llvm::GlobalVariable *CodeGenModule::CreateOrReplaceCXXRuntimeVariable(
@@ -4051,7 +3648,7 @@ llvm::GlobalVariable *CodeGenModule::CreateOrReplaceCXXRuntimeVariable(
 
   if (GV) {
     // Check if the variable has the right type.
-    if (GV->getValueType() == Ty)
+    if (GV->getType()->getElementType() == Ty)
       return GV;
 
     // Because C++ name mangling, the only way we can end up with an already
@@ -4100,10 +3697,11 @@ 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, Ty,
-                               getContext().getTargetAddressSpace(ASTTy), D,
-                               IsForDefinition);
+  return GetOrCreateLLVMGlobal(MangledName, PTy, D, IsForDefinition);
 }
 
 /// CreateRuntimeVariable - Create a new runtime global variable with the
@@ -4111,11 +3709,12 @@ llvm::Constant *CodeGenModule::GetAddrOfGlobalVar(const VarDecl *D,
 llvm::Constant *
 CodeGenModule::CreateRuntimeVariable(llvm::Type *Ty,
                                      StringRef Name) {
-  auto AddrSpace =
+  auto PtrTy =
       getContext().getLangOpts().OpenCL
-          ? getContext().getTargetAddressSpace(LangAS::opencl_global)
-          : 0;
-  auto *Ret = GetOrCreateLLVMGlobal(Name, Ty, AddrSpace, nullptr);
+          ? llvm::PointerType::get(
+                Ty, getContext().getTargetAddressSpace(LangAS::opencl_global))
+          : llvm::PointerType::getUnqual(Ty);
+  auto *Ret = GetOrCreateLLVMGlobal(Name, PtrTy, nullptr);
   setDSOLocal(cast<llvm::GlobalValue>(Ret->stripPointerCasts()));
   return Ret;
 }
@@ -4157,18 +3756,12 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
   if (LangOpts.OpenCL) {
     AddrSpace = D ? D->getType().getAddressSpace() : LangAS::opencl_global;
     assert(AddrSpace == LangAS::opencl_global ||
-           AddrSpace == LangAS::opencl_global_device ||
-           AddrSpace == LangAS::opencl_global_host ||
            AddrSpace == LangAS::opencl_constant ||
            AddrSpace == LangAS::opencl_local ||
            AddrSpace >= LangAS::FirstTargetAddressSpace);
     return AddrSpace;
   }
 
-  if (LangOpts.SYCLIsDevice &&
-      (!D || D->getType().getAddressSpace() == LangAS::Default))
-    return LangAS::sycl_global;
-
   if (LangOpts.CUDA && LangOpts.CUDAIsDevice) {
     if (D && D->hasAttr<CUDAConstantAttr>())
       return LangAS::cuda_constant;
@@ -4190,12 +3783,10 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
   return getTargetCodeGenInfo().getGlobalVarAddressSpace(*this, D);
 }
 
-LangAS CodeGenModule::GetGlobalConstantAddressSpace() const {
+LangAS CodeGenModule::getStringLiteralAddressSpace() 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;
@@ -4214,12 +3805,13 @@ castStringLiteralToDefaultAddressSpace(CodeGenModule &CGM,
                                        llvm::GlobalVariable *GV) {
   llvm::Constant *Cast = GV;
   if (!CGM.getLangOpts().OpenCL) {
-    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)));
+    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)));
+    }
   }
   return Cast;
 }
@@ -4309,7 +3901,7 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
       OpenMPRuntime->emitTargetGlobalVariable(D))
     return;
 
-  llvm::TrackingVH<llvm::Constant> Init;
+  llvm::Constant *Init = nullptr;
   bool NeedsGlobalCtor = false;
   bool NeedsGlobalDtor =
       D->needsDestruction(getContext()) == QualType::DK_cxx_destructor;
@@ -4326,20 +3918,17 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
       getLangOpts().CUDAIsDevice && D->hasAttr<CUDASharedAttr>();
   // Shadows of initialized device-side global variables are also left
   // undefined.
-  // Managed Variables should be initialized on both host side and device side.
   bool IsCUDAShadowVar =
-      !getLangOpts().CUDAIsDevice && !D->hasAttr<HIPManagedAttr>() &&
+      !getLangOpts().CUDAIsDevice &&
       (D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>() ||
        D->hasAttr<CUDASharedAttr>());
-  bool IsCUDADeviceShadowVar =
-      getLangOpts().CUDAIsDevice && !D->hasAttr<HIPManagedAttr>() &&
-      (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
-       D->getType()->isCUDADeviceBuiltinTextureType());
+  // HIP pinned shadow of initialized host-side global variables are also
+  // left undefined.
+  bool IsHIPPinnedShadowVar =
+      getLangOpts().CUDAIsDevice && D->hasAttr<HIPPinnedShadowAttr>();
   if (getLangOpts().CUDA &&
-      (IsCUDASharedVar || IsCUDAShadowVar || IsCUDADeviceShadowVar))
-    Init = llvm::UndefValue::get(getTypes().ConvertTypeForMem(ASTTy));
-  else if (D->hasAttr<LoaderUninitializedAttr>())
-    Init = llvm::UndefValue::get(getTypes().ConvertTypeForMem(ASTTy));
+      (IsCUDASharedVar || IsCUDAShadowVar || IsHIPPinnedShadowVar))
+    Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
   else if (!InitExpr) {
     // This is a tentative definition; tentative definitions are
     // implicitly initialized with { 0 }.
@@ -4355,8 +3944,9 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
   } else {
     initializedGlobalDecl = GlobalDecl(D);
     emitter.emplace(*this);
-    llvm::Constant *Initializer = emitter->tryEmitForInitializer(*InitDecl);
-    if (!Initializer) {
+    Init = emitter->tryEmitForInitializer(*InitDecl);
+
+    if (!Init) {
       QualType T = InitExpr->getType();
       if (D->getType()->isReferenceType())
         T = D->getType();
@@ -4369,7 +3959,6 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
         Init = llvm::UndefValue::get(getTypes().ConvertType(T));
       }
     } else {
-      Init = Initializer;
       // We don't need an initializer, so remove the entry for the delayed
       // initializer position (just in case this entry was delayed) if we
       // also don't need to register a destructor.
@@ -4397,7 +3986,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->getValueType() != InitType ||
+  if (!GV || GV->getType()->getElementType() != InitType ||
       GV->getType()->getAddressSpace() !=
           getContext().getTargetAddressSpace(GetGlobalVarAddressSpace(D))) {
 
@@ -4411,8 +4000,7 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
 
     // Replace all uses of the old global with the new global
     llvm::Constant *NewPtrForOldDecl =
-        llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(GV,
-                                                             Entry->getType());
+        llvm::ConstantExpr::getBitCast(GV, Entry->getType());
     Entry->replaceAllUsesWith(NewPtrForOldDecl);
 
     // Erase the old global, since it is no longer used.
@@ -4441,14 +4029,38 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
           (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()))
         GV->setExternallyInitialized(true);
     } else {
-      getCUDARuntime().internalizeDeviceSideVar(D, Linkage);
+      // Host-side shadows of external declarations of device-side
+      // global variables become internal definitions. These have to
+      // be internal in order to prevent name conflicts with global
+      // host variables with the same name in a different TUs.
+      if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
+          D->hasAttr<HIPPinnedShadowAttr>()) {
+        Linkage = llvm::GlobalValue::InternalLinkage;
+
+        // Shadow variables and their properties must be registered
+        // with CUDA runtime.
+        unsigned Flags = 0;
+        if (!D->hasDefinition())
+          Flags |= CGCUDARuntime::ExternDeviceVar;
+        if (D->hasAttr<CUDAConstantAttr>())
+          Flags |= CGCUDARuntime::ConstantDeviceVar;
+        // Extern global variables will be registered in the TU where they are
+        // defined.
+        if (!D->hasExternalStorage())
+          getCUDARuntime().registerDeviceVar(D, *GV, Flags);
+      } else if (D->hasAttr<CUDASharedAttr>())
+        // __shared__ variables are odd. Shadows do get created, but
+        // they are not registered with the CUDA runtime, so they
+        // can't really be used to access their device-side
+        // counterparts. It's not clear yet whether it's nvcc's bug or
+        // a feature, but we've got to do the same for compatibility.
+        Linkage = llvm::GlobalValue::InternalLinkage;
     }
-    getCUDARuntime().handleVarRegistration(D, *GV);
   }
 
-  GV->setInitializer(Init);
-  if (emitter)
-    emitter->finalize(GV);
+  if (!IsHIPPinnedShadowVar)
+    GV->setInitializer(Init);
+  if (emitter) emitter->finalize(GV);
 
   // If it is safe to mark the global 'constant', do so now.
   GV->setConstant(!NeedsGlobalCtor && !NeedsGlobalDtor &&
@@ -4463,24 +4075,17 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
 
   GV->setAlignment(getContext().getDeclAlign(D).getAsAlign());
 
-  // 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 &&
+  // 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 &&
       Context.getTargetInfo().getTriple().isOSDarwin() &&
-      !D->hasAttr<ConstInitAttr>())
+      !llvm::GlobalVariable::isLinkOnceLinkage(Linkage) &&
+      !llvm::GlobalVariable::isWeakLinkage(Linkage))
     Linkage = llvm::GlobalValue::InternalLinkage;
 
   GV->setLinkage(Linkage);
@@ -4529,8 +4134,9 @@ void CodeGenModule::EmitExternalVarDeclaration(const VarDecl *D) {
     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);
+      llvm::PointerType *PTy =
+          llvm::PointerType::get(Ty, getContext().getTargetAddressSpace(ASTTy));
+      llvm::Constant *GV = GetOrCreateLLVMGlobal(D->getName(), PTy, D);
       DI->EmitExternalVariable(
           cast<llvm::GlobalVariable>(GV->stripPointerCasts()), D);
     }
@@ -4653,16 +4259,13 @@ llvm::GlobalValue::LinkageTypes CodeGenModule::getLLVMLinkageForDeclarator(
   // and must all be equivalent. However, we are not allowed to
   // throw away these explicit instantiations.
   //
-  // CUDA/HIP: For -fno-gpu-rdc case, device code is limited to one TU,
+  // We don't currently support CUDA device code spread out across multiple TUs,
   // so say that CUDA templates are either external (for kernels) or internal.
-  // 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.
+  // This lets llvm perform aggressive inter-procedural optimizations.
   if (Linkage == GVA_StrongODR) {
-    if (getLangOpts().AppleKext)
+    if (Context.getLangOpts().AppleKext)
       return llvm::Function::ExternalLinkage;
-    if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
-        !getLangOpts().GPURelocatableDeviceCode)
+    if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice)
       return D->hasAttr<CUDAGlobalAttr>() ? llvm::Function::ExternalLinkage
                                           : llvm::Function::InternalLinkage;
     return llvm::Function::WeakODRLinkage;
@@ -4702,6 +4305,7 @@ static void replaceUsesOfNonProtoConstant(llvm::Constant *old,
 
   llvm::Type *newRetTy = newFn->getReturnType();
   SmallVector<llvm::Value*, 4> newArgs;
+  SmallVector<llvm::OperandBundleDef, 1> newBundles;
 
   for (llvm::Value::use_iterator ui = old->use_begin(), ue = old->use_end();
          ui != ue; ) {
@@ -4758,7 +4362,6 @@ static void replaceUsesOfNonProtoConstant(llvm::Constant *old,
     newArgs.append(callSite->arg_begin(), callSite->arg_begin() + argNo);
 
     // Copy over any operand bundles.
-    SmallVector<llvm::OperandBundleDef, 1> newBundles;
     callSite->getOperandBundlesAsDefs(newBundles);
 
     llvm::CallBase *newCall;
@@ -4825,6 +4428,11 @@ void CodeGenModule::HandleCXXStaticMemberVarInstantiation(VarDecl *VD) {
 
 void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
                                                  llvm::GlobalValue *GV) {
+  // Check if this must be emitted as declare variant.
+  if (LangOpts.OpenMP && OpenMPRuntime &&
+      OpenMPRuntime->emitDeclareVariant(GD, /*IsForDefinition=*/true))
+    return;
+
   const auto *D = cast<FunctionDecl>(GD.getDecl());
 
   // Compute the function info and LLVM type.
@@ -4832,7 +4440,7 @@ void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
   llvm::FunctionType *Ty = getTypes().GetFunctionType(FI);
 
   // Get or create the prototype for the function.
-  if (!GV || (GV->getValueType() != Ty))
+  if (!GV || (GV->getType()->getElementType() != Ty))
     GV = cast<llvm::GlobalValue>(GetAddrOfFunction(GD, Ty, /*ForVTable=*/false,
                                                    /*DontDefer=*/true,
                                                    ForDefinition));
@@ -4853,12 +4461,10 @@ void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
 
   MaybeHandleStaticInExternC(D, Fn);
 
-  maybeSetTrivialComdat(*D, *Fn);
 
-  // Set CodeGen attributes that represent floating point environment.
-  setLLVMFunctionFEnvAttributes(D, Fn);
+  maybeSetTrivialComdat(*D, *Fn);
 
-  CodeGenFunction(*this).GenerateCode(GD, Fn, FI);
+  CodeGenFunction(*this).GenerateCode(D, Fn, FI);
 
   setNonAliasAttributes(GD, Fn);
   SetLLVMFunctionAttributesForDefinition(D, Fn);
@@ -4866,7 +4472,7 @@ void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
   if (const ConstructorAttr *CA = D->getAttr<ConstructorAttr>())
     AddGlobalCtor(Fn, CA->getPriority());
   if (const DestructorAttr *DA = D->getAttr<DestructorAttr>())
-    AddGlobalDtor(Fn, DA->getPriority(), true);
+    AddGlobalDtor(Fn, DA->getPriority());
   if (D->hasAttr<AnnotateAttr>())
     AddGlobalAnnotations(D, Fn);
 }
@@ -4902,18 +4508,16 @@ void CodeGenModule::EmitAliasDefinition(GlobalDecl GD) {
                                       /*ForVTable=*/false);
     LT = getFunctionLinkage(GD);
   } else {
-    Aliasee = GetOrCreateLLVMGlobal(AA->getAliasee(), DeclTy, 0,
+    Aliasee = GetOrCreateLLVMGlobal(AA->getAliasee(),
+                                    llvm::PointerType::getUnqual(DeclTy),
                                     /*D=*/nullptr);
-    if (const auto *VD = dyn_cast<VarDecl>(GD.getDecl()))
-      LT = getLLVMLinkageVarDefinition(VD, D->getType().isConstQualified());
-    else
-      LT = getFunctionLinkage(GD);
+    LT = getLLVMLinkageVarDefinition(cast<VarDecl>(GD.getDecl()),
+                                     D->getType().isConstQualified());
   }
 
   // Create the new alias itself, but don't set a name yet.
-  unsigned AS = Aliasee->getType()->getPointerAddressSpace();
   auto *GA =
-      llvm::GlobalAlias::create(DeclTy, AS, LT, "", Aliasee, &getModule());
+      llvm::GlobalAlias::create(DeclTy, 0, LT, "", Aliasee, &getModule());
 
   if (Entry) {
     if (GA->getAliasee() == Entry) {
@@ -5118,7 +4722,7 @@ CodeGenModule::GetAddrOfConstantCFString(const StringLiteral *Literal) {
         DeclContext *DC = TranslationUnitDecl::castToDeclContext(TUDecl);
 
         const VarDecl *VD = nullptr;
-        for (const auto *Result : DC->lookup(&II))
+        for (const auto &Result : DC->lookup(&II))
           if ((VD = dyn_cast<VarDecl>(Result)))
             break;
 
@@ -5230,8 +4834,6 @@ 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:
@@ -5330,7 +4932,7 @@ GenerateStringLiteral(llvm::Constant *C, llvm::GlobalValue::LinkageTypes LT,
                       CodeGenModule &CGM, StringRef GlobalName,
                       CharUnits Alignment) {
   unsigned AddrSpace = CGM.getContext().getTargetAddressSpace(
-      CGM.GetGlobalConstantAddressSpace());
+      CGM.getStringLiteralAddressSpace());
 
   llvm::Module &M = CGM.getModule();
   // Create a global variable for this string
@@ -5457,21 +5059,8 @@ ConstantAddress CodeGenModule::GetAddrOfGlobalTemporary(
 
   CharUnits Align = getContext().getTypeAlignInChars(MaterializedType);
 
-  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);
-  }
+  if (llvm::Constant *Slot = MaterializedGlobalTemporaryMap[E])
+    return ConstantAddress(Slot, Align);
 
   // FIXME: If an externally-visible declaration extends multiple temporaries,
   // we need to give each temporary the same name in every translation unit (and
@@ -5550,17 +5139,7 @@ ConstantAddress CodeGenModule::GetAddrOfGlobalTemporary(
         *this, GV, AddrSpace, LangAS::Default,
         Type->getPointerTo(
             getContext().getTargetAddressSpace(LangAS::Default)));
-
-  // Update the map with the new temporary. If we created a placeholder above,
-  // replace it with the new global now.
-  llvm::Constant *&Entry = MaterializedGlobalTemporaryMap[E];
-  if (Entry) {
-    Entry->replaceAllUsesWith(
-        llvm::ConstantExpr::getBitCast(CV, Entry->getType()));
-    llvm::cast<llvm::GlobalVariable>(Entry)->eraseFromParent();
-  }
-  Entry = CV;
-
+  MaterializedGlobalTemporaryMap[E] = CV;
   return ConstantAddress(CV, Align);
 }
 
@@ -5686,11 +5265,6 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
   if (D->isTemplated())
     return;
 
-  // Consteval function shouldn't be emitted.
-  if (auto *FD = dyn_cast<FunctionDecl>(D))
-    if (FD->isConsteval())
-      return;
-
   switch (D->getKind()) {
   case Decl::CXXConversion:
   case Decl::CXXMethod:
@@ -5726,27 +5300,22 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
     break;
   case Decl::ClassTemplateSpecialization: {
     const auto *Spec = cast<ClassTemplateSpecializationDecl>(D);
-    if (CGDebugInfo *DI = getModuleDebugInfo())
-      if (Spec->getSpecializationKind() ==
-              TSK_ExplicitInstantiationDefinition &&
-          Spec->hasDefinition())
-        DI->completeTemplateDefinition(*Spec);
+    if (DebugInfo &&
+        Spec->getSpecializationKind() == TSK_ExplicitInstantiationDefinition &&
+        Spec->hasDefinition())
+      DebugInfo->completeTemplateDefinition(*Spec);
   } LLVM_FALLTHROUGH;
-  case Decl::CXXRecord: {
-    CXXRecordDecl *CRD = cast<CXXRecordDecl>(D);
-    if (CGDebugInfo *DI = getModuleDebugInfo()) {
-      if (CRD->hasDefinition())
-        DI->EmitAndRetainType(getContext().getRecordType(cast<RecordDecl>(D)));
+  case Decl::CXXRecord:
+    if (DebugInfo) {
       if (auto *ES = D->getASTContext().getExternalSource())
         if (ES->hasExternalDefinitions(D) == ExternalASTSource::EK_Never)
-          DI->completeUnusedClass(*CRD);
+          DebugInfo->completeUnusedClass(cast<CXXRecordDecl>(*D));
     }
     // Emit any static data members, they may be definitions.
-    for (auto *I : CRD->decls())
+    for (auto *I : cast<CXXRecordDecl>(D)->decls())
       if (isa<VarDecl>(I) || isa<CXXRecordDecl>(I))
         EmitTopLevelDecl(I);
     break;
-  }
     // No code generation needed.
   case Decl::UsingShadow:
   case Decl::ClassTemplate:
@@ -5762,19 +5331,15 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
   case Decl::Using:          // using X; [C++]
     if (CGDebugInfo *DI = getModuleDebugInfo())
         DI->EmitUsingDecl(cast<UsingDecl>(*D));
-    break;
-  case Decl::UsingEnum: // using enum X; [C++]
-    if (CGDebugInfo *DI = getModuleDebugInfo())
-      DI->EmitUsingEnumDecl(cast<UsingEnumDecl>(*D));
-    break;
+    return;
   case Decl::NamespaceAlias:
     if (CGDebugInfo *DI = getModuleDebugInfo())
         DI->EmitNamespaceAlias(cast<NamespaceAliasDecl>(*D));
-    break;
+    return;
   case Decl::UsingDirective: // using namespace X; [C++]
     if (CGDebugInfo *DI = getModuleDebugInfo())
       DI->EmitUsingDirective(cast<UsingDirectiveDecl>(*D));
-    break;
+    return;
   case Decl::CXXConstructor:
     getCXXABI().EmitCXXConstructors(cast<CXXConstructorDecl>(D));
     break;
@@ -5865,9 +5430,6 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
     // File-scope asm is ignored during device-side OpenMP compilation.
     if (LangOpts.OpenMPIsDevice)
       break;
-    // File-scope asm is ignored during device-side SYCL compilation.
-    if (LangOpts.SYCLIsDevice)
-      break;
     auto *AD = cast<FileScopeAsmDecl>(D);
     getModule().appendModuleInlineAsm(AD->getAsmString()->getString());
     break;
@@ -5925,7 +5487,6 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
     break;
 
   case Decl::OMPAllocate:
-    EmitOMPAllocateDecl(cast<OMPAllocateDecl>(D));
     break;
 
   case Decl::OMPDeclareReduction:
@@ -5940,25 +5501,6 @@ void CodeGenModule::EmitTopLevelDecl(Decl *D) {
     EmitOMPRequiresDecl(cast<OMPRequiresDecl>(D));
     break;
 
-  case Decl::Typedef:
-  case Decl::TypeAlias: // using foo = bar; [C++11]
-    if (CGDebugInfo *DI = getModuleDebugInfo())
-      DI->EmitAndRetainType(
-          getContext().getTypedefType(cast<TypedefNameDecl>(D)));
-    break;
-
-  case Decl::Record:
-    if (CGDebugInfo *DI = getModuleDebugInfo())
-      if (cast<RecordDecl>(D)->getDefinition())
-        DI->EmitAndRetainType(getContext().getRecordType(cast<RecordDecl>(D)));
-    break;
-
-  case Decl::Enum:
-    if (CGDebugInfo *DI = getModuleDebugInfo())
-      if (cast<EnumDecl>(D)->getDefinition())
-        DI->EmitAndRetainType(getContext().getEnumType(cast<EnumDecl>(D)));
-    break;
-
   default:
     // Make sure we handled everything we should, every other kind is a
     // non-top-level decl.  FIXME: Would be nice to have an isTopLevelDeclKind
@@ -5980,10 +5522,10 @@ void CodeGenModule::AddDeferredUnusedCoverageMapping(Decl *D) {
   case Decl::CXXConstructor:
   case Decl::CXXDestructor: {
     if (!cast<FunctionDecl>(D)->doesThisDeclarationHaveABody())
-      break;
+      return;
     SourceManager &SM = getContext().getSourceManager();
     if (LimitedCoverage && SM.getMainFileID() != SM.getFileID(D->getBeginLoc()))
-      break;
+      return;
     auto I = DeferredEmptyCoverageMappingDecls.find(D);
     if (I == DeferredEmptyCoverageMappingDecls.end())
       DeferredEmptyCoverageMappingDecls[D] = true;
@@ -6049,17 +5591,6 @@ 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) {
@@ -6095,7 +5626,7 @@ void CodeGenModule::EmitStaticExternCAliases() {
     IdentifierInfo *Name = I.first;
     llvm::GlobalValue *Val = I.second;
     if (Val && !getModule().getNamedValue(Name->getName()))
-      addCompilerUsedGlobal(llvm::GlobalAlias::create(Name->getName(), Val));
+      addUsedGlobal(llvm::GlobalAlias::create(Name->getName(), Val));
   }
 }
 
@@ -6174,6 +5705,21 @@ 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())
@@ -6196,14 +5742,39 @@ 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 ||
-      (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
-       getTriple().isNVPTX()))
+  if ((!ForEH && !getLangOpts().RTTI) || getLangOpts().CUDAIsDevice)
     return llvm::Constant::getNullValue(Int8PtrTy);
 
   if (ForEH && Ty->isObjCObjectPointerType() &&
@@ -6337,116 +5908,13 @@ llvm::SanitizerStatReport &CodeGenModule::getSanStats() {
 
   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);
-  auto *Call = CGF.EmitRuntimeCall(
-      CreateRuntimeFunction(FTy, "__translate_sampler_initializer"), {C});
-  return Call;
-}
-
-CharUnits CodeGenModule::getNaturalPointeeTypeAlignment(
-    QualType T, LValueBaseInfo *BaseInfo, TBAAAccessInfo *TBAAInfo) {
-  return getNaturalTypeAlignment(T->getPointeeType(), BaseInfo, TBAAInfo,
-                                 /* forPointeeType= */ true);
-}
-
-CharUnits CodeGenModule::getNaturalTypeAlignment(QualType T,
-                                                 LValueBaseInfo *BaseInfo,
-                                                 TBAAAccessInfo *TBAAInfo,
-                                                 bool forPointeeType) {
-  if (TBAAInfo)
-    *TBAAInfo = getTBAAAccessInfo(T);
-
-  // FIXME: This duplicates logic in ASTContext::getTypeAlignIfKnown. But
-  // that doesn't return the information we need to compute BaseInfo.
-
-  // Honor alignment typedef attributes even on incomplete types.
-  // We also honor them straight for C++ class types, even as pointees;
-  // there's an expressivity gap here.
-  if (auto TT = T->getAs<TypedefType>()) {
-    if (auto Align = TT->getDecl()->getMaxAlignment()) {
-      if (BaseInfo)
-        *BaseInfo = LValueBaseInfo(AlignmentSource::AttributedType);
-      return getContext().toCharUnitsFromBits(Align);
-    }
-  }
-
-  bool AlignForArray = T->isArrayType();
-
-  // Analyze the base element type, so we don't get confused by incomplete
-  // array types.
-  T = getContext().getBaseElementType(T);
-
-  if (T->isIncompleteType()) {
-    // We could try to replicate the logic from
-    // ASTContext::getTypeAlignIfKnown, but nothing uses the alignment if the
-    // type is incomplete, so it's impossible to test. We could try to reuse
-    // getTypeAlignIfKnown, but that doesn't return the information we need
-    // to set BaseInfo.  So just ignore the possibility that the alignment is
-    // greater than one.
-    if (BaseInfo)
-      *BaseInfo = LValueBaseInfo(AlignmentSource::Type);
-    return CharUnits::One();
-  }
-
-  if (BaseInfo)
-    *BaseInfo = LValueBaseInfo(AlignmentSource::Type);
-
-  CharUnits Alignment;
-  const CXXRecordDecl *RD;
-  if (T.getQualifiers().hasUnaligned()) {
-    Alignment = CharUnits::One();
-  } else if (forPointeeType && !AlignForArray &&
-             (RD = T->getAsCXXRecordDecl())) {
-    // For C++ class pointees, we don't know whether we're pointing at a
-    // base or a complete object, so we generally need to use the
-    // non-virtual alignment.
-    Alignment = getClassPointerAlignment(RD);
-  } else {
-    Alignment = getContext().getTypeAlignInChars(T);
-  }
-
-  // Cap to the global maximum type alignment unless the alignment
-  // was somehow explicit on the type.
-  if (unsigned MaxAlign = getLangOpts().MaxTypeAlign) {
-    if (Alignment.getQuantity() > MaxAlign &&
-        !getContext().isAlignmentRequired(T))
-      Alignment = CharUnits::fromQuantity(MaxAlign);
-  }
-  return Alignment;
-}
-
-bool CodeGenModule::stopAutoInit() {
-  unsigned StopAfter = getContext().getLangOpts().TrivialAutoVarInitStopAfter;
-  if (StopAfter) {
-    // This number is positive only when -ftrivial-auto-var-init-stop-after=* is
-    // used
-    if (NumAutoVarInit >= StopAfter) {
-      return true;
-    }
-    if (!NumAutoVarInit) {
-      unsigned DiagID = getDiags().getCustomDiagID(
-          DiagnosticsEngine::Warning,
-          "-ftrivial-auto-var-init-stop-after=%0 has been enabled to limit the "
-          "number of times ftrivial-auto-var-init=%1 gets applied.");
-      getDiags().Report(DiagID)
-          << StopAfter
-          << (getContext().getLangOpts().getTrivialAutoVarInit() ==
-                      LangOptions::TrivialAutoVarInitKind::Zero
-                  ? "zero"
-                  : "pattern");
-    }
-    ++NumAutoVarInit;
-  }
-  return false;
-}
-
-void CodeGenModule::printPostfixForExternalizedStaticVar(
-    llvm::raw_ostream &OS) const {
-  OS << ".static." << getContext().getCUIDHash();
+  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});
 }