[Clang][CodeGen]`vtable`, `typeinfo` et al. are globals

[AlexVlx]

Summary

All data structures and values associated with handling virtual functions / inheritance, as well as RTTI, are globals and thus can only reside in the global address space. This was not taken fully taken into account because for most targets, global & generic appear to coincide. However, on targets where global & generic ASes differ (e.g. AMDGPU), this was problematic, since it led to the generation of invalid bitcasts (which would trigger asserts in Debug) and less than optimal code. This patch does two things:

• ensures that vtables, vptrs, vtts, typeinfo are generated in the right AS, and populated accordingly;
• removes a bunch of bitcasts which look like left-overs from the typed ptr era.

This is a bit more noisy than I'd have liked, but functionality is somewhat spread out. There's one bit of less than ideal code, stemming from the fact that functions are in the generic AS, and thus it's necessary to insert a constexpr cast from generic to global when populating the vtable. Adjusting appears disruptive enough to prefer to do it separately (unless I missed something obvious).

[AlexVlx]

Rebased.

[AlexVlx]

clang-format

[AlexVlx]

Gentle ping.

[AlexVlx]

Fixed issue found via internal testing (thanks @yaxunl).

[yaxunl]

In D153092#4447445, @AlexVlx wrote:

Fixed issue found via internal testing (thanks @yaxunl).

Can we add a test to cover the regression found via internal testing? Thanks.

[AlexVlx]

Add missing test for vtable initializers on the __device__ side.

[AlexVlx]

In D153092#4448514, @yaxunl wrote:

In D153092#4447445, @AlexVlx wrote:

Fixed issue found via internal testing (thanks @yaxunl).

Can we add a test to cover the regression found via internal testing? Thanks.

Done, I had forgotten about that.

[yaxunl]

This could be a good chance to switch VT to constant address space instead of global address space. AFAIK if a target has global addr space they usually also has constant addr space since they usually support OpenCL or CUDA/HIP. Is there any reason we cannot introduce a CGM.ConstantGlobalsInt8PtrTy and use it for VT instead?

[AlexVlx]

In D153092#4452070, @yaxunl wrote:

This could be a good chance to switch VT to constant address space instead of global address space. AFAIK if a target has global addr space they usually also has constant addr space since they usually support OpenCL or CUDA/HIP. Is there any reason we cannot introduce a CGM.ConstantGlobalsInt8PtrTy and use it for VT instead?

I did give this some thought and the benefits are somewhat unclear to the point of being ultimately counterproductive. Note that these are already marked constant, which IIRC is / was going to be enough to get most of the benefits, at least on our back-end. Furthermore, the semantics of the constant address space are a bit weird in something like OpenCL e.g. A pointer that points to the constant address space cannot be cast or implicitly converted to the generic address space.. This would lead to weirdness when composing with CUDA / HIP, where constant is treated as device, which is to say global. IIRC, you are also meant to use magical interfaces to write into constant from the host, which a loader wouldn't necessarily do. Overall, I think that the OCL formulation of constant is actually meant to allow for relatively strange things like loading things into ROM or having different pointer types (be it width or canonicity). TL;DR, I am concerned that a target could validly have the constant addr space be disjoint from generic/flat, with no viable way to even cast between the two. We could say “yes, but this is not that constant”, but then if OCL ever starts supporting dynamic polymorphism it would get confusing.

[yaxunl]

In D153092#4452251, @AlexVlx wrote:

In D153092#4452070, @yaxunl wrote:

This could be a good chance to switch VT to constant address space instead of global address space. AFAIK if a target has global addr space they usually also has constant addr space since they usually support OpenCL or CUDA/HIP. Is there any reason we cannot introduce a CGM.ConstantGlobalsInt8PtrTy and use it for VT instead?

I did give this some thought and the benefits are somewhat unclear to the point of being ultimately counterproductive. Note that these are already marked constant, which IIRC is / was going to be enough to get most of the benefits, at least on our back-end. Furthermore, the semantics of the constant address space are a bit weird in something like OpenCL e.g. A pointer that points to the constant address space cannot be cast or implicitly converted to the generic address space.. This would lead to weirdness when composing with CUDA / HIP, where constant is treated as device, which is to say global. IIRC, you are also meant to use magical interfaces to write into constant from the host, which a loader wouldn't necessarily do. Overall, I think that the OCL formulation of constant is actually meant to allow for relatively strange things like loading things into ROM or having different pointer types (be it width or canonicity). TL;DR, I am concerned that a target could validly have the constant addr space be disjoint from generic/flat, with no viable way to even cast between the two. We could say “yes, but this is not that constant”, but then if OCL ever starts supporting dynamic polymorphism it would get confusing.

I agree that constant address space may be target dependent about what can be put there. I also agree there is

In D153092#4452251, @AlexVlx wrote:

In D153092#4452070, @yaxunl wrote:

This could be a good chance to switch VT to constant address space instead of global address space. AFAIK if a target has global addr space they usually also has constant addr space since they usually support OpenCL or CUDA/HIP. Is there any reason we cannot introduce a CGM.ConstantGlobalsInt8PtrTy and use it for VT instead?

I did give this some thought and the benefits are somewhat unclear to the point of being ultimately counterproductive. Note that these are already marked constant, which IIRC is / was going to be enough to get most of the benefits, at least on our back-end. Furthermore, the semantics of the constant address space are a bit weird in something like OpenCL e.g. A pointer that points to the constant address space cannot be cast or implicitly converted to the generic address space.. This would lead to weirdness when composing with CUDA / HIP, where constant is treated as device, which is to say global. IIRC, you are also meant to use magical interfaces to write into constant from the host, which a loader wouldn't necessarily do. Overall, I think that the OCL formulation of constant is actually meant to allow for relatively strange things like loading things into ROM or having different pointer types (be it width or canonicity). TL;DR, I am concerned that a target could validly have the constant addr space be disjoint from generic/flat, with no viable way to even cast between the two. We could say “yes, but this is not that constant”, but then if OCL ever starts supporting dynamic polymorphism it would get confusing.

Sounds reasonable.

LGTM. Thanks.

[AlexVlx]

Thank you @yaxunl. @rjmccall / @efriedma any input on this? I'd like to try landing it next week to unblock some additional work. Thanks!

[AlexVlx]

Rebase.

[AlexVlx]

Ping.

[AlexVlx]

Rebase.

[efriedma]

clang/lib/CodeGen/CGVTables.cpp

836

If I follow correctly, the fnPtr is guaranteed to be in the global address-space, but GetAddrOfFunction returns a generic pointer. So the vtable entries are in the global address-space for efficiency? Seems reasonable.

There isn't really much point to explicitly checking FnAS != GVAS before the call to getAddrSpaceCast; getAddrSpaceCast does the same check internally.

[AlexVlx]

clang/lib/CodeGen/CGVTables.cpp

836

Right, we know that these are going to be in global memory, and there is overhead when dealing with flat/generic. I would've liked to remove the check, but I think that's infeasible with how getAddrSpaceCast is currently implemented, because it asserts on castIsValid; addrspacecasts from the same AS to the same AS are invalid, so the assert flares on targets where FnAS == GVAS (e.g. x86). This is not super ergonomic IMHO, as it should be valid & a NOP just returning the source, but that's a change that would be required to allow deleting the silly check, I believe.

[bjope]

clang/lib/CodeGen/CGVTables.cpp

693

I noticed that we have some old fixes downstream that conflicts with the changes you've made here. I thought that perhaps we could get rid of those now when you've fixed the code upstream.

Isn't the VTable holding function pointers when not using the relative layout, and then this should be a pointer to the ProgramAddressSpace and not a pointer to the DefaultGlobalsAddressSpace?

Downstream we've been using a special FnVoidPtrTy here. Defined as FnVoidPtrTy = Int8Ty->getPointerTo(DL.getProgramAddressSpace());.

[rjmccall]

clang/lib/CodeGen/CGVTables.cpp

693

It's a mix. The type_info pointer should be in the global address space (although it would be forgivable to just use the default address space), the top, vbase, and vcall offsets are all ptrdiff_ts (presumably the same size as the default address space), and the virtual functions are function pointers. If we're going to support a target where those can be different sizes, we probably need to start computing a byte layout of the v-table and doing byte GEPs into it, because our current IR patterns are naively assuming the components are all the same size.

clang/lib/CodeGen/CGVTT.cpp

return CGVT.GenerateConstructionVTable(MostDerivedClass,
                                          AddressPoints);
 }
 
 void
 CodeGenVTables::EmitVTTDefinition(llvm::GlobalVariable *VTT,
                                   llvm::GlobalVariable::LinkageTypes Linkage,
                                   const CXXRecordDecl *RD) {
   VTTBuilder Builder(CGM.getContext(), RD, /*GenerateDefinition=*/true);
-  llvm::ArrayType *ArrayType =
-      llvm::ArrayType::get(CGM.Int8PtrTy, Builder.getVTTComponents().size());
+  llvm::ArrayType *ArrayType = llvm::ArrayType::get(
+      CGM.GlobalsInt8PtrTy, Builder.getVTTComponents().size());
 
   SmallVector<llvm::GlobalVariable *, 8> VTables;
   SmallVector<VTableAddressPointsMapTy, 8> VTableAddressPoints;
   for (const VTTVTable *i = Builder.getVTTVTables().begin(),
                        *e = Builder.getVTTVTables().end(); i != e; ++i) {
     VTableAddressPoints.push_back(VTableAddressPointsMapTy());
     VTables.push_back(GetAddrOfVTTVTable(*this, CGM, RD, *i, Linkage,
                                          VTableAddressPoints.back()));
llvm::Value *Idxs[] = {
        llvm::ConstantInt::get(CGM.Int32Ty, AddressPoint.VTableIndex),
        llvm::ConstantInt::get(CGM.Int32Ty, AddressPoint.AddressPointIndex),
      };
 
      llvm::Constant *Init = llvm::ConstantExpr::getGetElementPtr(
          VTable->getValueType(), VTable, Idxs, /*InBounds=*/true,
          /*InRangeIndex=*/1);
 
-     Init = llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(Init,
-                                                                 CGM.Int8PtrTy);
-
      VTTComponents.push_back(Init);
   }
 
   llvm::Constant *Init = llvm::ConstantArray::get(ArrayType, VTTComponents);
 
   VTT->setInitializer(Init);
 
   // Set the correct linkage.
cast<ItaniumMangleContext>(CGM.getCXXABI().getMangleContext())
       .mangleCXXVTT(RD, Out);
   StringRef Name = OutName.str();
 
   // This will also defer the definition of the VTT.
   (void) CGM.getCXXABI().getAddrOfVTable(RD, CharUnits());
 
   VTTBuilder Builder(CGM.getContext(), RD, /*GenerateDefinition=*/false);
 
-  llvm::ArrayType *ArrayType =
-    llvm::ArrayType::get(CGM.Int8PtrTy, Builder.getVTTComponents().size());
-  llvm::Align Align = CGM.getDataLayout().getABITypeAlign(CGM.Int8PtrTy);
+  llvm::ArrayType *ArrayType = llvm::ArrayType::get(
+      CGM.GlobalsInt8PtrTy, Builder.getVTTComponents().size());
+  llvm::Align Align = CGM.getDataLayout().getABITypeAlign(CGM.GlobalsInt8PtrTy);
 
   llvm::GlobalVariable *GV = CGM.CreateOrReplaceCXXRuntimeVariable(
       Name, ArrayType, llvm::GlobalValue::ExternalLinkage, Align);
   GV->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
   CGM.setGVProperties(GV, RD);
   return GV;
 }
 

clang/lib/CodeGen/CGVTables.cpp

 
 bool CodeGenVTables::useRelativeLayout() const {
   return UseRelativeLayout(CGM);
 }
 
 llvm::Type *CodeGenModule::getVTableComponentType() const {
   if (UseRelativeLayout(*this))
     return Int32Ty;
-  return Int8PtrTy;
+  return GlobalsInt8PtrTy;
 }
 
 llvm::Type *CodeGenVTables::getVTableComponentType() const {
   return CGM.getVTableComponentType();
 }
 
 static void AddPointerLayoutOffset(const CodeGenModule &CGM,
                                    ConstantArrayBuilder &builder,
                                    CharUnits offset) {
   builder.add(llvm::ConstantExpr::getIntToPtr(
       llvm::ConstantInt::get(CGM.PtrDiffTy, offset.getQuantity()),
-      CGM.Int8PtrTy));
+      CGM.GlobalsInt8PtrTy));
 }
 
 static void AddRelativeLayoutOffset(const CodeGenModule &CGM,
                                     ConstantArrayBuilder &builder,
                                     CharUnits offset) {
   builder.add(llvm::ConstantInt::get(CGM.Int32Ty, offset.getQuantity()));
 }
 
case VTableComponent::CK_OffsetToTop:
     return addOffsetConstant(CGM, builder, component.getOffsetToTop());
 
   case VTableComponent::CK_RTTI:
     if (useRelativeLayout())
       return addRelativeComponent(builder, rtti, vtableAddressPoint,
                                   vtableHasLocalLinkage,
                                   /*isCompleteDtor=*/false);
     else
-      return builder.add(llvm::ConstantExpr::getBitCast(rtti, CGM.Int8PtrTy));
+      return builder.add(rtti);
 
   case VTableComponent::CK_FunctionPointer:
   case VTableComponent::CK_CompleteDtorPointer:
   case VTableComponent::CK_DeletingDtorPointer: {
     GlobalDecl GD = component.getGlobalDecl();
 
     if (CGM.getLangOpts().CUDA) {
       // Emit NULL for methods we can't codegen on this
       // side. Otherwise we'd end up with vtable with unresolved
       // references.
       const CXXMethodDecl *MD = cast<CXXMethodDecl>(GD.getDecl());
       // OK on device side: functions w/ __device__ attribute
       // OK on host side: anything except __device__-only functions.
       bool CanEmitMethod =
           CGM.getLangOpts().CUDAIsDevice
               ? MD->hasAttr<CUDADeviceAttr>()
               : (MD->hasAttr<CUDAHostAttr>() || !MD->hasAttr<CUDADeviceAttr>());
       if (!CanEmitMethod)
-        return builder.add(llvm::ConstantExpr::getNullValue(CGM.Int8PtrTy));
+        return builder.add(
+            llvm::ConstantExpr::getNullValue(CGM.GlobalsInt8PtrTy));
       // Method is acceptable, continue processing as usual.
     }
 
     auto getSpecialVirtualFn = [&](StringRef name) -> llvm::Constant * {
       // FIXME(PR43094): When merging comdat groups, lld can select a local
       // symbol as the signature symbol even though it cannot be accessed
       // outside that symbol's TU. The relative vtables ABI would make
       // __cxa_pure_virtual and __cxa_deleted_virtual local symbols, and
       // depending on link order, the comdat groups could resolve to the one
       // with the local symbol. As a temporary solution, fill these components
       // with zero. We shouldn't be calling these in the first place anyway.
       if (useRelativeLayout())
-        return llvm::ConstantPointerNull::get(CGM.Int8PtrTy);
+        return llvm::ConstantPointerNull::get(CGM.GlobalsInt8PtrTy);
 
       // For NVPTX devices in OpenMP emit special functon as null pointers,
       // otherwise linking ends up with unresolved references.
       if (CGM.getLangOpts().OpenMP && CGM.getLangOpts().OpenMPIsTargetDevice &&
           CGM.getTriple().isNVPTX())
-        return llvm::ConstantPointerNull::get(CGM.Int8PtrTy);
+        return llvm::ConstantPointerNull::get(CGM.GlobalsInt8PtrTy);
       llvm::FunctionType *fnTy =
           llvm::FunctionType::get(CGM.VoidTy, /*isVarArg=*/false);
       llvm::Constant *fn = cast<llvm::Constant>(
           CGM.CreateRuntimeFunction(fnTy, name).getCallee());
       if (auto f = dyn_cast<llvm::Function>(fn))
         f->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
-      return llvm::ConstantExpr::getBitCast(fn, CGM.Int8PtrTy);
+      return fn;
     };
 
     llvm::Constant *fnPtr;
 
     // Pure virtual member functions.
     if (cast<CXXMethodDecl>(GD.getDecl())->isPure()) {
       if (!PureVirtualFn)
         PureVirtualFn =
// Otherwise we can use the method definition directly.
       llvm::Type *fnTy = CGM.getTypes().GetFunctionTypeForVTable(GD);
       fnPtr = CGM.GetAddrOfFunction(GD, fnTy, /*ForVTable=*/true);
     }
 
     if (useRelativeLayout()) {
       return addRelativeComponent(
           builder, fnPtr, vtableAddressPoint, vtableHasLocalLinkage,
           component.getKind() == VTableComponent::CK_CompleteDtorPointer);
-    } else
-      return builder.add(llvm::ConstantExpr::getBitCast(fnPtr, CGM.Int8PtrTy));
+    } else {
+      // TODO: this icky and only exists due to functions being in the generic
+      //       address space, rather than the global one, even though they are
+      //       globals;  fixing said issue might be intrusive, and will be done
+      //       later.
+      unsigned FnAS = fnPtr->getType()->getPointerAddressSpace();
+      unsigned GVAS = CGM.GlobalsInt8PtrTy->getPointerAddressSpace();
+
+      if (FnAS != GVAS)
+        fnPtr =
+            llvm::ConstantExpr::getAddrSpaceCast(fnPtr, CGM.GlobalsInt8PtrTy);
+      return builder.add(fnPtr);
+    }
   }
 
   case VTableComponent::CK_UnusedFunctionPointer:
     if (useRelativeLayout())
       return builder.add(llvm::ConstantExpr::getNullValue(CGM.Int32Ty));
     else
-      return builder.addNullPointer(CGM.Int8PtrTy);
+      return builder.addNullPointer(CGM.GlobalsInt8PtrTy);
   }
 
   llvm_unreachable("Unexpected vtable component kind");
 }
 
 llvm::Type *CodeGenVTables::getVTableType(const VTableLayout &layout) {
   SmallVector<llvm::Type *, 4> tys;
   llvm::Type *componentType = getVTableComponentType();

clang/lib/CodeGen/CodeGenModule.cpp

 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().OpenMPIsTargetDevice &&
        getTriple().isNVPTX()))
-    return llvm::Constant::getNullValue(Int8PtrTy);
+    return llvm::Constant::getNullValue(GlobalsInt8PtrTy);
 
   if (ForEH && Ty->isObjCObjectPointerType() &&
       LangOpts.ObjCRuntime.isGNUFamily())
     return ObjCRuntime->GetEHType(Ty);
 
   return getCXXABI().getAddrOfRTTIDescriptor(Ty);
 }
 

clang/lib/CodeGen/ItaniumCXXABI.cpp

CGCallee ItaniumCXXABI::EmitLoadOfMemberFunctionPointer(
   Builder.CreateCondBr(IsVirtual, FnVirtual, FnNonVirtual);
 
   // In the virtual path, the adjustment left 'This' pointing to the
   // vtable of the correct base subobject.  The "function pointer" is an
   // offset within the vtable (+1 for the virtual flag on non-ARM).
   CGF.EmitBlock(FnVirtual);
 
   // Cast the adjusted this to a pointer to vtable pointer and load.
-  llvm::Type *VTableTy = Builder.getInt8PtrTy();
+  llvm::Type *VTableTy = CGF.CGM.GlobalsInt8PtrTy;
   CharUnits VTablePtrAlign =
     CGF.CGM.getDynamicOffsetAlignment(ThisAddr.getAlignment(), RD,
                                       CGF.getPointerAlign());
   llvm::Value *VTable = CGF.GetVTablePtr(
       Address(This, ThisAddr.getElementType(), VTablePtrAlign), VTableTy, RD);
 
   // Apply the offset.
   // On ARM64, to reserve extra space in virtual member function pointers,
llvm::Value *ItaniumCXXABI::getVTableAddressPointInStructorWithVTT(
 
   // Get the secondary vpointer index.
   uint64_t VirtualPointerIndex =
       CGM.getVTables().getSecondaryVirtualPointerIndex(VTableClass, Base);
 
   /// Load the VTT.
   llvm::Value *VTT = CGF.LoadCXXVTT();
   if (VirtualPointerIndex)
-    VTT = CGF.Builder.CreateConstInBoundsGEP1_64(
-        CGF.VoidPtrTy, VTT, VirtualPointerIndex);
+    VTT = CGF.Builder.CreateConstInBoundsGEP1_64(CGF.GlobalsVoidPtrTy, VTT,
+                                                 VirtualPointerIndex);
 
   // And load the address point from the VTT.
-  return CGF.Builder.CreateAlignedLoad(CGF.VoidPtrTy, VTT,
+  return CGF.Builder.CreateAlignedLoad(CGF.GlobalsVoidPtrTy, VTT,
                                        CGF.getPointerAlign());
 }
 
 llvm::Constant *ItaniumCXXABI::getVTableAddressPointForConstExpr(
     BaseSubobject Base, const CXXRecordDecl *VTableClass) {
   return getVTableAddressPoint(Base, VTableClass);
 }
 
llvm::GlobalVariable *ItaniumCXXABI::getAddrOfVTable(const CXXRecordDecl *RD,
   SmallString<256> Name;
   llvm::raw_svector_ostream Out(Name);
   getMangleContext().mangleCXXVTable(RD, Out);
 
   const VTableLayout &VTLayout =
       CGM.getItaniumVTableContext().getVTableLayout(RD);
   llvm::Type *VTableType = CGM.getVTables().getVTableType(VTLayout);
 
-  // Use pointer alignment for the vtable. Otherwise we would align them based
-  // on the size of the initializer which doesn't make sense as only single
-  // values are read.
+  // Use pointer to global alignment for the vtable. Otherwise we would align
+  // them based on the size of the initializer which doesn't make sense as only
+  // single values are read.
+  LangAS AS = CGM.GetGlobalVarAddressSpace(nullptr);
   unsigned PAlign = CGM.getItaniumVTableContext().isRelativeLayout()
                         ? 32
-                        : CGM.getTarget().getPointerAlign(LangAS::Default);
+                        : CGM.getTarget().getPointerAlign(AS);
 
   VTable = CGM.CreateOrReplaceCXXRuntimeVariable(
       Name, VTableType, llvm::GlobalValue::ExternalLinkage,
       getContext().toCharUnitsFromBits(PAlign).getAsAlign());
   VTable->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
 
   // In MS C++ if you have a class with virtual functions in which you are using
   // selective member import/export, then all virtual functions must be exported
llvm::GlobalVariable *ItaniumCXXABI::getAddrOfVTable(const CXXRecordDecl *RD,
   return VTable;
 }
 
 CGCallee ItaniumCXXABI::getVirtualFunctionPointer(CodeGenFunction &CGF,
                                                   GlobalDecl GD,
                                                   Address This,
                                                   llvm::Type *Ty,
                                                   SourceLocation Loc) {
-  llvm::Type *PtrTy = llvm::PointerType::getUnqual(CGF.getLLVMContext());
+  llvm::Type *PtrTy = CGM.GlobalsInt8PtrTy;
   auto *MethodDecl = cast<CXXMethodDecl>(GD.getDecl());
   llvm::Value *VTable = CGF.GetVTablePtr(This, PtrTy, MethodDecl->getParent());
 
   uint64_t VTableIndex = CGM.getItaniumVTableContext().getMethodVTableIndex(GD);
   llvm::Value *VFunc;
   if (CGF.ShouldEmitVTableTypeCheckedLoad(MethodDecl->getParent())) {
     VFunc = CGF.EmitVTableTypeCheckedLoad(
         MethodDecl->getParent(), VTable, PtrTy,
ItaniumRTTIBuilder::GetAddrOfExternalRTTIDescriptor(QualType Ty) {
   // Look for an existing global.
   llvm::GlobalVariable *GV = CGM.getModule().getNamedGlobal(Name);
 
   if (!GV) {
     // Create a new global variable.
     // Note for the future: If we would ever like to do deferred emission of
     // RTTI, check if emitting vtables opportunistically need any adjustment.
 
-    GV = new llvm::GlobalVariable(CGM.getModule(), CGM.Int8PtrTy,
-                                  /*isConstant=*/true,
-                                  llvm::GlobalValue::ExternalLinkage, nullptr,
-                                  Name);
+    GV = new llvm::GlobalVariable(
+        CGM.getModule(), CGM.GlobalsInt8PtrTy,
+        /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage, nullptr, Name);
     const CXXRecordDecl *RD = Ty->getAsCXXRecordDecl();
     CGM.setGVProperties(GV, RD);
     // Import the typeinfo symbol when all non-inline virtual methods are
     // imported.
     if (CGM.getTarget().hasPS4DLLImportExport()) {
       if (RD && CXXRecordAllNonInlineVirtualsHaveAttr<DLLImportAttr>(RD)) {
         GV->setDLLStorageClass(llvm::GlobalVariable::DLLImportStorageClass);
         CGM.setDSOLocal(GV);
       }
     }
   }
 
-  return llvm::ConstantExpr::getBitCast(GV, CGM.Int8PtrTy);
+  return GV;
 }
 
 /// TypeInfoIsInStandardLibrary - Given a builtin type, returns whether the type
 /// info for that type is defined in the standard library.
 static bool TypeInfoIsInStandardLibrary(const BuiltinType *Ty) {
   // Itanium C++ ABI 2.9.2:
   //   Basic type information (e.g. for "int", "bool", etc.) will be kept in
   //   the run-time support library. Specifically, the run-time support
#include "clang/AST/TypeNodes.inc"
   }
 
   llvm::Constant *VTable = nullptr;
 
   // Check if the alias exists. If it doesn't, then get or create the global.
   if (CGM.getItaniumVTableContext().isRelativeLayout())
     VTable = CGM.getModule().getNamedAlias(VTableName);
   if (!VTable)
-    VTable = CGM.getModule().getOrInsertGlobal(VTableName, CGM.Int8PtrTy);
+    VTable =
+        CGM.getModule().getOrInsertGlobal(VTableName, CGM.GlobalsInt8PtrTy);
 
   CGM.setDSOLocal(cast<llvm::GlobalValue>(VTable->stripPointerCasts()));
 
   llvm::Type *PtrDiffTy =
       CGM.getTypes().ConvertType(CGM.getContext().getPointerDiffType());
 
   // The vtable address point is 2.
   if (CGM.getItaniumVTableContext().isRelativeLayout()) {
     // The vtable address point is 8 bytes after its start:
     // 4 for the offset to top + 4 for the relative offset to rtti.
     llvm::Constant *Eight = llvm::ConstantInt::get(CGM.Int32Ty, 8);
-    VTable = llvm::ConstantExpr::getBitCast(VTable, CGM.Int8PtrTy);
     VTable =
         llvm::ConstantExpr::getInBoundsGetElementPtr(CGM.Int8Ty, VTable, Eight);
   } else {
     llvm::Constant *Two = llvm::ConstantInt::get(PtrDiffTy, 2);
-    VTable = llvm::ConstantExpr::getInBoundsGetElementPtr(CGM.Int8PtrTy, VTable,
-                                                          Two);
+    VTable = llvm::ConstantExpr::getInBoundsGetElementPtr(CGM.GlobalsInt8PtrTy,
+                                                          VTable, Two);
   }
-  VTable = llvm::ConstantExpr::getBitCast(VTable, CGM.Int8PtrTy);
 
   Fields.push_back(VTable);
 }
 
 /// Return the linkage that the type info and type info name constants
 /// should have for the given type.
 static llvm::GlobalVariable::LinkageTypes getTypeInfoLinkage(CodeGenModule &CGM,
                                                              QualType Ty) {
llvm::Constant *ItaniumRTTIBuilder::BuildTypeInfo(QualType Ty) {
   llvm::raw_svector_ostream Out(Name);
   CGM.getCXXABI().getMangleContext().mangleCXXRTTI(Ty, Out);
 
   llvm::GlobalVariable *OldGV = CGM.getModule().getNamedGlobal(Name);
   if (OldGV && !OldGV->isDeclaration()) {
     assert(!OldGV->hasAvailableExternallyLinkage() &&
            "available_externally typeinfos not yet implemented");
 
-    return llvm::ConstantExpr::getBitCast(OldGV, CGM.Int8PtrTy);
+    return OldGV;
   }
 
   // Check if there is already an external RTTI descriptor for this type.
   if (IsStandardLibraryRTTIDescriptor(Ty) ||
       ShouldUseExternalRTTIDescriptor(CGM, Ty))
     return GetAddrOfExternalRTTIDescriptor(Ty);
 
   // Emit the standard library with external linkage.
llvm::Constant *ItaniumRTTIBuilder::BuildTypeInfo(
   if (RTTIUniqueness != ItaniumCXXABI::RUK_Unique) {
     // The flag is the sign bit, which on ARM64 is defined to be clear
     // for global pointers.  This is very ARM64-specific.
     TypeNameField = llvm::ConstantExpr::getPtrToInt(TypeName, CGM.Int64Ty);
     llvm::Constant *flag =
         llvm::ConstantInt::get(CGM.Int64Ty, ((uint64_t)1) << 63);
     TypeNameField = llvm::ConstantExpr::getAdd(TypeNameField, flag);
     TypeNameField =
-        llvm::ConstantExpr::getIntToPtr(TypeNameField, CGM.Int8PtrTy);
+        llvm::ConstantExpr::getIntToPtr(TypeNameField, CGM.GlobalsInt8PtrTy);
   } else {
-    TypeNameField = llvm::ConstantExpr::getBitCast(TypeName, CGM.Int8PtrTy);
+    TypeNameField = TypeName;
   }
   Fields.push_back(TypeNameField);
 
   switch (Ty->getTypeClass()) {
 #define TYPE(Class, Base)
 #define ABSTRACT_TYPE(Class, Base)
 #define NON_CANONICAL_UNLESS_DEPENDENT_TYPE(Class, Base) case Type::Class:
 #define NON_CANONICAL_TYPE(Class, Base) case Type::Class:
if (OldGV) {
     OldGV->replaceAllUsesWith(NewPtr);
     OldGV->eraseFromParent();
   }
 
   if (CGM.supportsCOMDAT() && GV->isWeakForLinker())
     GV->setComdat(M.getOrInsertComdat(GV->getName()));
 
   CharUnits Align = CGM.getContext().toCharUnitsFromBits(
-      CGM.getTarget().getPointerAlign(LangAS::Default));
+      CGM.getTarget().getPointerAlign(CGM.GetGlobalVarAddressSpace(nullptr)));
   GV->setAlignment(Align.getAsAlign());
 
   // The Itanium ABI specifies that type_info objects must be globally
   // unique, with one exception: if the type is an incomplete class
   // type or a (possibly indirect) pointer to one.  That exception
   // affects the general case of comparing type_info objects produced
   // by the typeid operator, which is why the comparison operators on
   // std::type_info generally use the type_info name pointers instead
#include "clang/AST/TypeNodes.inc"
   TypeName->setDLLStorageClass(DLLStorageClass);
   GV->setDLLStorageClass(CGM.getTarget().hasPS4DLLImportExport()
                              ? GVDLLStorageClass
                              : DLLStorageClass);
 
   TypeName->setPartition(CGM.getCodeGenOpts().SymbolPartition);
   GV->setPartition(CGM.getCodeGenOpts().SymbolPartition);
 
-  return llvm::ConstantExpr::getBitCast(GV, CGM.Int8PtrTy);
+  return GV;
 }
 
 /// BuildObjCObjectTypeInfo - Build the appropriate kind of type_info
 /// for the given Objective-C object type.
 void ItaniumRTTIBuilder::BuildObjCObjectTypeInfo(const ObjCObjectType *OT) {
   // Drop qualifiers.
   const Type *T = OT->getBaseType().getTypePtr();
   assert(isa<BuiltinType>(T) || isa<ObjCInterfaceType>(T));

clang/test/CodeGenCXX/vtable-align-address-space.cpp

+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -std=c++11 -emit-llvm -o - | FileCheck %s
+
+struct A {
+  virtual void f();
+  virtual void g();
+  virtual void h();
+};
+
+void A::f() {}
+
+// CHECK: @_ZTV1A ={{.*}} unnamed_addr addrspace(1) constant { [5 x ptr addrspace(1)] } { [5 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) @_ZTI1A, ptr addrspace(1) addrspacecast (ptr @_ZN1A1fEv to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @_ZN1A1gEv to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @_ZN1A1hEv to ptr addrspace(1))]
+// CHECK: @_ZTS1A ={{.*}} constant [3 x i8] c"1A\00", align 1
+// CHECK: @_ZTI1A ={{.*}} addrspace(1) constant { ptr addrspace(1), ptr addrspace(1) } { ptr addrspace(1) getelementptr inbounds (ptr addrspace(1), ptr addrspace(1) @_ZTVN10__cxxabiv117__class_type_infoE, i64 2), ptr addrspace(1) @_ZTS1A }, align 8

clang/test/CodeGenCXX/vtable-assume-load-address-space.cpp

+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -std=c++11 -emit-llvm -o %t.ll -O1 -disable-llvm-passes -fms-extensions -fstrict-vtable-pointers
+// FIXME: Assume load should not require -fstrict-vtable-pointers
+
+// RUN: FileCheck --check-prefix=CHECK1 --input-file=%t.ll %s
+// RUN: FileCheck --check-prefix=CHECK2 --input-file=%t.ll %s
+// RUN: FileCheck --check-prefix=CHECK3 --input-file=%t.ll %s
+// RUN: FileCheck --check-prefix=CHECK4 --input-file=%t.ll %s
+// RUN: FileCheck --check-prefix=CHECK5 --input-file=%t.ll %s
+// RUN: FileCheck --check-prefix=CHECK6 --input-file=%t.ll %s
+// RUN: FileCheck --check-prefix=CHECK7 --input-file=%t.ll %s
+// RUN: FileCheck --check-prefix=CHECK8 --input-file=%t.ll %s
+namespace test1 {
+
+struct A {
+  A();
+  virtual void foo();
+};
+
+struct B : A {
+  virtual void foo();
+};
+
+void g(A *a) { a->foo(); }
+
+// CHECK1-LABEL: define{{.*}} void @_ZN5test14fooAEv()
+// CHECK1: call void @_ZN5test11AC1Ev(ptr
+// CHECK1: %[[VTABLE:.*]] = load ptr addrspace(1), ptr %{{.*}}
+// CHECK1: %[[CMP:.*]] = icmp eq ptr addrspace(1) %[[VTABLE]], getelementptr inbounds ({ [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test11AE, i32 0, inrange i32 0, i32 2)
+// CHECK1: call void @llvm.assume(i1 %[[CMP]])
+// CHECK1-LABEL: {{^}}}
+
+void fooA() {
+  A a;
+  g(&a);
+}
+
+// CHECK1-LABEL: define{{.*}} void @_ZN5test14fooBEv()
+// CHECK1: call void @_ZN5test11BC1Ev(ptr {{[^,]*}} %{{.*}})
+// CHECK1: %[[VTABLE:.*]] = load ptr addrspace(1), ptr %{{.*}}
+// CHECK1: %[[CMP:.*]] = icmp eq ptr addrspace(1) %[[VTABLE]], getelementptr inbounds ({ [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test11BE, i32 0, inrange i32 0, i32 2)
+// CHECK1: call void @llvm.assume(i1 %[[CMP]])
+// CHECK1-LABEL: {{^}}}
+
+void fooB() {
+  B b;
+  g(&b);
+}
+// there should not be any assumes in the ctor that calls base ctor
+// CHECK1-LABEL: define linkonce_odr void @_ZN5test11BC2Ev(ptr
+// CHECK1-NOT: @llvm.assume(
+// CHECK1-LABEL: {{^}}}
+}
+namespace test2 {
+struct A {
+  A();
+  virtual void foo();
+};
+
+struct B {
+  B();
+  virtual void bar();
+};
+
+struct C : A, B {
+  C();
+  virtual void foo();
+};
+void g(A *a) { a->foo(); }
+void h(B *b) { b->bar(); }
+
+// CHECK2-LABEL: define{{.*}} void @_ZN5test24testEv()
+// CHECK2: call void @_ZN5test21CC1Ev(ptr
+// CHECK2: %[[VTABLE:.*]] = load ptr addrspace(1), ptr {{.*}}
+// CHECK2: %[[CMP:.*]] = icmp eq ptr addrspace(1) %[[VTABLE]], getelementptr inbounds ({ [3 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test21CE, i32 0, inrange i32 0, i32 2)
+// CHECK2: call void @llvm.assume(i1 %[[CMP]])
+
+// CHECK2: %[[ADD_PTR:.*]] = getelementptr inbounds i8, ptr %{{.*}}, i64 8
+// CHECK2: %[[VTABLE2:.*]] = load ptr addrspace(1), ptr %[[ADD_PTR]]
+// CHECK2: %[[CMP2:.*]] = icmp eq ptr addrspace(1) %[[VTABLE2]], getelementptr inbounds ({ [3 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test21CE, i32 0, inrange i32 1, i32 2)
+// CHECK2: call void @llvm.assume(i1 %[[CMP2]])
+
+// CHECK2: call void @_ZN5test21gEPNS_1AE(
+// CHECK2-LABEL: {{^}}}
+
+void test() {
+  C c;
+  g(&c);
+  h(&c);
+}
+}
+
+namespace test3 {
+struct A {
+  A();
+};
+
+struct B : A {
+  B();
+  virtual void foo();
+};
+
+struct C : virtual A, B {
+  C();
+  virtual void foo();
+};
+void g(B *a) { a->foo(); }
+
+// CHECK3-LABEL: define{{.*}} void @_ZN5test34testEv()
+// CHECK3: call void @_ZN5test31CC1Ev(ptr
+// CHECK3: %[[CMP:.*]] = icmp eq ptr addrspace(1) %{{.*}}, getelementptr inbounds ({ [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test31CE, i32 0, inrange i32 0, i32 3)
+// CHECK3: call void @llvm.assume(i1 %[[CMP]])
+// CHECK3-LABLEL: }
+void test() {
+  C c;
+  g(&c);
+}
+} // test3
+
+namespace test4 {
+struct A {
+  A();
+  virtual void foo();
+};
+
+struct B : virtual A {
+  B();
+  virtual void foo();
+};
+struct C : B {
+  C();
+  virtual void foo();
+};
+
+void g(C *c) { c->foo(); }
+
+// CHECK4-LABEL: define{{.*}} void @_ZN5test44testEv()
+// CHECK4: call void @_ZN5test41CC1Ev(ptr
+// CHECK4: %[[VTABLE:.*]] = load ptr addrspace(1), ptr %{{.*}}
+// CHECK4: %[[CMP:.*]] = icmp eq ptr addrspace(1) %[[VTABLE]], getelementptr inbounds ({ [5 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test41CE, i32 0, inrange i32 0, i32 4)
+// CHECK4: call void @llvm.assume(i1 %[[CMP]]
+
+// CHECK4: %[[VTABLE2:.*]] = load ptr addrspace(1), ptr %{{.*}}
+// CHECK4: %[[CMP2:.*]] = icmp eq ptr addrspace(1) %[[VTABLE2]], getelementptr inbounds ({ [5 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5test41CE, i32 0, inrange i32 0, i32 4)
+// CHECK4: call void @llvm.assume(i1 %[[CMP2]])
+// CHECK4-LABEL: {{^}}}
+
+void test() {
+  C c;
+  g(&c);
+}
+} // test4
+
+namespace test6 {
+struct A {
+  A();
+  virtual void foo();
+  virtual ~A() {}
+};
+struct B : A {
+  B();
+};
+// FIXME: Because A's vtable is external, and no virtual functions are hidden,
+// it's safe to generate assumption loads.
+// CHECK5-LABEL: define{{.*}} void @_ZN5test61gEv()
+// CHECK5: call void @_ZN5test61AC1Ev(
+// CHECK5-NOT: call void @llvm.assume(
+
+// We can't emit assumption loads for B, because if we would refer to vtable
+// it would refer to functions that will not be able to find (like implicit
+// inline destructor).
+
+// CHECK5-LABEL:   call void @_ZN5test61BC1Ev(
+// CHECK5-NOT: call void @llvm.assume(
+// CHECK5-LABEL: {{^}}}
+void g() {
+  A *a = new A;
+  B *b = new B;
+}
+}
+
+namespace test7 {
+// Because A's key function is defined here, vtable is generated in this TU
+// CHECK6: @_ZTVN5test71AE ={{.*}} unnamed_addr addrspace(1) constant
+struct A {
+  A();
+  virtual void foo();
+  virtual void bar();
+};
+void A::foo() {}
+
+// CHECK6-LABEL: define{{.*}} void @_ZN5test71gEv()
+// CHECK6: call void @_ZN5test71AC1Ev(
+// CHECK6: call void @llvm.assume(
+// CHECK6-LABEL: {{^}}}
+void g() {
+  A *a = new A();
+  a->bar();
+}
+}
+
+namespace test8 {
+
+struct A {
+  virtual void foo();
+  virtual void bar();
+};
+
+// CHECK7-DAG: @_ZTVN5test81BE = available_externally unnamed_addr addrspace(1) constant
+struct B : A {
+  B();
+  void foo();
+  void bar();
+};
+
+// CHECK7-DAG: @_ZTVN5test81CE = linkonce_odr unnamed_addr addrspace(1) constant
+struct C : A {
+  C();
+  void bar();
+  void foo() {}
+};
+inline void C::bar() {}
+
+struct D : A {
+  D();
+  void foo();
+  void inline bar();
+};
+void D::bar() {}
+
+// CHECK7-DAG: @_ZTVN5test81EE = linkonce_odr unnamed_addr addrspace(1) constant
+struct E : A {
+  E();
+};
+
+// CHECK7-LABEL: define{{.*}} void @_ZN5test81bEv()
+// CHECK7: call void @llvm.assume(
+// CHECK7-LABEL: {{^}}}
+void b() {
+  B b;
+  b.bar();
+}
+
+// FIXME: C has inline virtual functions which prohibits as from generating
+// assumption loads, but because vtable is generated in this TU (key function
+// defined here) it would be correct to refer to it.
+// CHECK7-LABEL: define{{.*}} void @_ZN5test81cEv()
+// CHECK7-NOT: call void @llvm.assume(
+// CHECK7-LABEL: {{^}}}
+void c() {
+  C c;
+  c.bar();
+}
+
+// FIXME: We could generate assumption loads here.
+// CHECK7-LABEL: define{{.*}} void @_ZN5test81dEv()
+// CHECK7-NOT: call void @llvm.assume(
+// CHECK7-LABEL: {{^}}}
+void d() {
+  D d;
+  d.bar();
+}
+
+// CHECK7-LABEL: define{{.*}} void @_ZN5test81eEv()
+// CHECK7: call void @llvm.assume(
+// CHECK7-LABEL: {{^}}}
+void e() {
+  E e;
+  e.bar();
+}
+}
+
+namespace test9 {
+
+struct S {
+  S();
+  __attribute__((visibility("hidden"))) virtual void doStuff();
+};
+
+// CHECK8-LABEL: define{{.*}} void @_ZN5test94testEv()
+// CHECK8-NOT: @llvm.assume(
+// CHECK8: }
+void test() {
+  S *s = new S();
+  s->doStuff();
+  delete s;
+}
+}
+

clang/test/CodeGenCXX/vtable-consteval-address-space.cpp

+// RUN: %clang_cc1 -std=c++20 -triple=amdgcn-amd-amdhsa %s -emit-llvm -o - | FileCheck %s --check-prefix=ITANIUM --implicit-check-not=DoNotEmit
+
+// FIXME: The MSVC ABI rule in use here was discussed with MS folks prior to
+// them implementing virtual consteval functions, but we do not know for sure
+// if this is the ABI rule they will use.
+
+// ITANIUM-DAG: @_ZTV1A = {{.*}} addrspace(1) constant { [2 x ptr addrspace(1)] } {{.*}} null, {{.*}} @_ZTI1A
+struct A {
+  virtual consteval void DoNotEmit_f() {}
+};
+// ITANIUM-DAG: @a = addrspace(1) global { {{.*}} ptr addrspace(1) @_ZTV1A,
+A a;
+
+// ITANIUM-DAG: @_ZTV1B = {{.*}} addrspace(1) constant { [4 x ptr addrspace(1)] } {{.*}} addrspace(1) null, ptr addrspace(1) @_ZTI1B, ptr addrspace(1) addrspacecast (ptr @_ZN1B1fEv to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @_ZN1B1hEv to ptr addrspace(1))
+struct B {
+  virtual void f() {}
+  virtual consteval void DoNotEmit_g() {}
+  virtual void h() {}
+};
+// ITANIUM-DAG: @b = addrspace(1) global { {{.*}} @_ZTV1B,
+B b;
+
+// ITANIUM-DAG: @_ZTV1C = {{.*}} addrspace(1) constant { [4 x ptr addrspace(1)] } {{.*}} addrspace(1) null, ptr addrspace(1) @_ZTI1C, ptr addrspace(1) addrspacecast (ptr @_ZN1CD1Ev to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @_ZN1CD0Ev to ptr addrspace(1))
+struct C {
+  virtual ~C() = default;
+  virtual consteval C &operator=(const C&) = default;
+};
+// ITANIUM-DAG: @c = addrspace(1) global { {{.*}} @_ZTV1C,
+C c;
+
+// ITANIUM-DAG: @_ZTV1D = {{.*}} addrspace(1) constant { [4 x ptr addrspace(1)] } {{.*}} addrspace(1) null, ptr addrspace(1) @_ZTI1D, ptr addrspace(1) addrspacecast (ptr @_ZN1DD1Ev to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @_ZN1DD0Ev to ptr addrspace(1))
+struct D : C {};
+// ITANIUM-DAG: @d = addrspace(1) global { ptr addrspace(1) } { {{.*}} @_ZTV1D,
+D d;
+
+// ITANIUM-DAG: @_ZTV1E = {{.*}} addrspace(1) constant { [3 x ptr addrspace(1)] } {{.*}} addrspace(1) null, ptr addrspace(1) @_ZTI1E, ptr addrspace(1) addrspacecast (ptr @_ZN1E1fEv to ptr addrspace(1))
+struct E { virtual void f() {} };
+// ITANIUM-DAG: @e = addrspace(1) global { {{.*}} @_ZTV1E,
+E e;
+
+// ITANIUM-DAG: @_ZTV1F = {{.*}} addrspace(1) constant { [3 x ptr addrspace(1)] } {{.*}} addrspace(1) null, ptr addrspace(1) @_ZTI1F, ptr addrspace(1) addrspacecast (ptr @_ZN1E1fEv to ptr addrspace(1))
+struct F : E { virtual consteval void DoNotEmit_g(); };
+// ITANIUM-DAG: @f = addrspace(1) global { ptr addrspace(1) } { {{.*}} @_ZTV1F,
+F f;

clang/test/CodeGenCXX/vtable-constexpr-address-space.cpp

+// RUN: %clang_cc1 -std=c++20 -triple=amdgcn-amd-amdhsa %s -emit-llvm -o - | FileCheck %s --implicit-check-not=DoNotEmit
+
+// constexpr virtual functions can be called at runtime and go in the vtable as
+// normal. But they are implicitly inline so are never the key function.
+
+struct DoNotEmit {
+  virtual constexpr void f();
+};
+constexpr void DoNotEmit::f() {}
+
+// CHECK-DAG: @_ZTV1B = {{.*}} addrspace(1) constant { [3 x ptr addrspace(1)] } { {{.*}} null, {{.*}} @_ZTI1B, {{.*}} @_ZN1B1fEv
+struct B {
+  // CHECK-DAG: define {{.*}} @_ZN1B1fEv
+  virtual constexpr void f() {}
+};
+B b;
+
+struct CBase {
+  virtual constexpr void f(); // not key function
+};
+
+// CHECK-DAG: @_ZTV1C = {{.*}} addrspace(1) constant {{.*}} null, {{.*}} @_ZTI1C, {{.*}} @_ZN1C1fEv
+struct C : CBase {
+  void f(); // key function
+};
+// CHECK-DAG: define {{.*}} @_ZN1C1fEv
+void C::f() {}

clang/test/CodeGenCXX/vtable-key-function-address-space.cpp

+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck %s
+// PR5697
+namespace PR5697 {
+struct A {
+  virtual void f() { }
+  A();
+  A(int);
+};
+
+// A does not have a key function, so the first constructor we emit should
+// cause the vtable to be defined (without assertions.)
+// CHECK: @_ZTVN6PR56971AE = linkonce_odr unnamed_addr addrspace(1) constant
+A::A() { }
+A::A(int) { }
+}
+
+// Make sure that we don't assert when building the vtable for a class
+// template specialization or explicit instantiation with a key
+// function.
+template<typename T>
+struct Base {
+  virtual ~Base();
+};
+
+template<typename T>
+struct Derived : public Base<T> { };
+
+template<>
+struct Derived<char> : public Base<char> {
+  virtual void anchor();
+};
+
+void Derived<char>::anchor() { }

clang/test/CodeGenCXX/vtable-layout-extreme-address-space.cpp

+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -std=c++11 -emit-llvm-only -fdump-vtable-layouts 2>&1 | FileCheck %s
+
+// A collection of big class hierarchies and their vtables.
+
+namespace Test1 {
+
+class C0
+{
+};
+class C1
+ :  virtual public C0
+{
+  int k0;
+};
+class C2
+ :  public C0
+ ,  virtual public C1
+{
+  int k0;
+};
+class C3
+ :  virtual public C0
+ ,  virtual public C1
+ ,  public C2
+{
+  int k0;
+  int k1;
+  int k2;
+  int k3;
+};
+class C4
+ :  public C2
+ ,  virtual public C3
+ ,  public C0
+{
+  int k0;
+};
+class C5
+ :  public C0
+ ,  virtual public C4
+ ,  public C2
+ ,  public C1
+ ,  virtual public C3
+{
+  int k0;
+};
+class C6
+ :  virtual public C3
+ ,  public C0
+ ,  public C5
+ ,  public C4
+ ,  public C1
+{
+  int k0;
+};
+class C7
+ :  virtual public C5
+ ,  virtual public C6
+ ,  virtual public C3
+ ,  public C4
+ ,  virtual public C2
+{
+  int k0;
+  int k1;
+};
+class C8
+ :  public C7
+ ,  public C5
+ ,  public C3
+ ,  virtual public C4
+ ,  public C1
+ ,  public C2
+{
+  int k0;
+  int k1;
+};
+
+// CHECK:     Vtable for 'Test1::C9' (87 entries).
+// CHECK-NEXT:   0 | vbase_offset (344)
+// CHECK-NEXT:   1 | vbase_offset (312)
+// CHECK-NEXT:   2 | vbase_offset (184)
+// CHECK-NEXT:   3 | vbase_offset (168)
+// CHECK-NEXT:   4 | vbase_offset (120)
+// CHECK-NEXT:   5 | vbase_offset (48)
+// CHECK-NEXT:   6 | vbase_offset (148)
+// CHECK-NEXT:   7 | vbase_offset (152)
+// CHECK-NEXT:   8 | offset_to_top (0)
+// CHECK-NEXT:   9 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C2, 0) vtable address --
+// CHECK-NEXT:       -- (Test1::C9, 0) vtable address --
+// CHECK-NEXT:  10 | void Test1::C9::f()
+// CHECK-NEXT:  11 | vbase_offset (104)
+// CHECK-NEXT:  12 | vbase_offset (132)
+// CHECK-NEXT:  13 | vbase_offset (136)
+// CHECK-NEXT:  14 | offset_to_top (-16)
+// CHECK-NEXT:  15 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C2, 16) vtable address --
+// CHECK-NEXT:       -- (Test1::C4, 16) vtable address --
+// CHECK-NEXT:  16 | vbase_offset (72)
+// CHECK-NEXT:  17 | vbase_offset (120)
+// CHECK-NEXT:  18 | vbase_offset (100)
+// CHECK-NEXT:  19 | vbase_offset (104)
+// CHECK-NEXT:  20 | offset_to_top (-48)
+// CHECK-NEXT:  21 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C2, 48) vtable address --
+// CHECK-NEXT:       -- (Test1::C5, 48) vtable address --
+// CHECK-NEXT:       -- (Test1::C6, 48) vtable address --
+// CHECK-NEXT:  22 | vbase_offset (84)
+// CHECK-NEXT:  23 | offset_to_top (-64)
+// CHECK-NEXT:  24 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C1, 64) vtable address --
+// CHECK-NEXT:  25 | vbase_offset (32)
+// CHECK-NEXT:  26 | vbase_offset (60)
+// CHECK-NEXT:  27 | vbase_offset (64)
+// CHECK-NEXT:  28 | offset_to_top (-88)
+// CHECK-NEXT:  29 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C2, 88) vtable address --
+// CHECK-NEXT:       -- (Test1::C4, 88) vtable address --
+// CHECK-NEXT:  30 | vbase_offset (44)
+// CHECK-NEXT:  31 | offset_to_top (-104)
+// CHECK-NEXT:  32 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C1, 104) vtable address --
+// CHECK-NEXT:  33 | vbase_offset (28)
+// CHECK-NEXT:  34 | vbase_offset (32)
+// CHECK-NEXT:  35 | offset_to_top (-120)
+// CHECK-NEXT:  36 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C2, 120) vtable address --
+// CHECK-NEXT:       -- (Test1::C3, 120) vtable address --
+// CHECK-NEXT:  37 | vbase_offset (-4)
+// CHECK-NEXT:  38 | offset_to_top (-152)
+// CHECK-NEXT:  39 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C1, 152) vtable address --
+// CHECK-NEXT:  40 | vbase_offset (-48)
+// CHECK-NEXT:  41 | vbase_offset (-20)
+// CHECK-NEXT:  42 | vbase_offset (-16)
+// CHECK-NEXT:  43 | offset_to_top (-168)
+// CHECK-NEXT:  44 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C2, 168) vtable address --
+// CHECK-NEXT:       -- (Test1::C4, 168) vtable address --
+// CHECK-NEXT:  45 | vbase_offset (160)
+// CHECK-NEXT:  46 | vbase_offset (-136)
+// CHECK-NEXT:  47 | vbase_offset (-16)
+// CHECK-NEXT:  48 | vbase_offset (128)
+// CHECK-NEXT:  49 | vbase_offset (-64)
+// CHECK-NEXT:  50 | vbase_offset (-36)
+// CHECK-NEXT:  51 | vbase_offset (-32)
+// CHECK-NEXT:  52 | offset_to_top (-184)
+// CHECK-NEXT:  53 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C2, 184) vtable address --
+// CHECK-NEXT:       -- (Test1::C4, 184) vtable address --
+// CHECK-NEXT:       -- (Test1::C7, 184) vtable address --
+// CHECK-NEXT:       -- (Test1::C8, 184) vtable address --
+// CHECK-NEXT:  54 | vbase_offset (-88)
+// CHECK-NEXT:  55 | vbase_offset (-40)
+// CHECK-NEXT:  56 | vbase_offset (-60)
+// CHECK-NEXT:  57 | vbase_offset (-56)
+// CHECK-NEXT:  58 | offset_to_top (-208)
+// CHECK-NEXT:  59 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C2, 208) vtable address --
+// CHECK-NEXT:       -- (Test1::C5, 208) vtable address --
+// CHECK-NEXT:  60 | vbase_offset (-76)
+// CHECK-NEXT:  61 | offset_to_top (-224)
+// CHECK-NEXT:  62 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C1, 224) vtable address --
+// CHECK-NEXT:  63 | vbase_offset (-92)
+// CHECK-NEXT:  64 | vbase_offset (-88)
+// CHECK-NEXT:  65 | offset_to_top (-240)
+// CHECK-NEXT:  66 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C2, 240) vtable address --
+// CHECK-NEXT:       -- (Test1::C3, 240) vtable address --
+// CHECK-NEXT:  67 | vbase_offset (-124)
+// CHECK-NEXT:  68 | offset_to_top (-272)
+// CHECK-NEXT:  69 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C1, 272) vtable address --
+// CHECK-NEXT:  70 | vbase_offset (-140)
+// CHECK-NEXT:  71 | vbase_offset (-136)
+// CHECK-NEXT:  72 | offset_to_top (-288)
+// CHECK-NEXT:  73 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C2, 288) vtable address --
+// CHECK-NEXT:  74 | vbase_offset (-192)
+// CHECK-NEXT:  75 | vbase_offset (-144)
+// CHECK-NEXT:  76 | vbase_offset (-164)
+// CHECK-NEXT:  77 | vbase_offset (-160)
+// CHECK-NEXT:  78 | offset_to_top (-312)
+// CHECK-NEXT:  79 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C2, 312) vtable address --
+// CHECK-NEXT:       -- (Test1::C5, 312) vtable address --
+// CHECK-NEXT:  80 | vbase_offset (-180)
+// CHECK-NEXT:  81 | offset_to_top (-328)
+// CHECK-NEXT:  82 | Test1::C9 RTTI
+// CHECK-NEXT:       -- (Test1::C1, 328) vtable address --
+// CHECK-NEXT:  83 | vbase_offset (-196)
+// CHECK-NEXT:  84 | vbase_offset (-192)
+// CHECK-NEXT:  85 | offset_to_top (-344)
+// CHECK-NEXT:  86 | Test1::C9 RTTI
+class C9
+ :  virtual public C6
+ ,  public C2
+ ,  public C4
+ ,  virtual public C8
+{
+  int k0;
+  int k1;
+  int k2;
+  int k3;
+  virtual void f();
+};
+void C9::f() { }
+
+}

clang/test/CodeGenCXX/vtable-linkage-address-space.cpp

+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -emit-llvm -o %t
+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -emit-llvm -std=c++03 -o %t.03
+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -emit-llvm -std=c++11 -o %t.11
+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -disable-llvm-passes -O3 -emit-llvm -o %t.opt
+// RUN: FileCheck %s < %t
+// RUN: FileCheck %s < %t.03
+// RUN: FileCheck %s < %t.11
+// RUN: FileCheck --check-prefix=CHECK-OPT %s < %t.opt
+
+namespace {
+  struct A {
+    virtual void f() { }
+  };
+}
+
+void f() { A b; }
+
+struct B {
+  B();
+  virtual void f();
+};
+
+B::B() { }
+
+struct C : virtual B {
+  C();
+  virtual void f() { }
+};
+
+C::C() { }
+
+struct D {
+  virtual void f();
+};
+
+void D::f() { }
+
+static struct : D { } e;
+
+// Force 'e' to be constructed and therefore have a vtable defined.
+void use_e() {
+  e.f();
+}
+
+// The destructor is the key function.
+template<typename T>
+struct E {
+  virtual ~E();
+};
+
+template<typename T> E<T>::~E() { }
+
+// Anchor is the key function
+template<>
+struct E<char> {
+  virtual void anchor();
+};
+
+void E<char>::anchor() { }
+
+template struct E<short>;
+extern template struct E<int>;
+
+void use_E() {
+  E<int> ei;
+  (void)ei;
+  E<long> el;
+  (void)el;
+}
+
+// No key function
+template<typename T>
+struct F {
+  virtual void foo() { }
+};
+
+// No key function
+template<>
+struct F<char> {
+  virtual void foo() { }
+};
+
+template struct F<short>;
+extern template struct F<int>;
+
+void use_F() {
+  F<char> fc;
+  fc.foo();
+  F<int> fi;
+  fi.foo();
+  F<long> fl;
+  (void)fl;
+}
+
+// B has a key function that is not defined in this translation unit so its vtable
+// has external linkage.
+// CHECK-DAG: @_ZTV1B = external unnamed_addr addrspace(1) constant
+
+// C has no key function, so its vtable should have weak_odr linkage
+// and hidden visibility (rdar://problem/7523229).
+// CHECK-DAG: @_ZTV1C = linkonce_odr unnamed_addr addrspace(1) constant {{.*}}, comdat, align 8{{$}}
+// CHECK-DAG: @_ZTS1C = linkonce_odr addrspace(1) constant {{.*}}, comdat, align 1{{$}}
+// CHECK-DAG: @_ZTI1C = linkonce_odr addrspace(1) constant {{.*}}, comdat, align 8{{$}}
+// CHECK-DAG: @_ZTT1C = linkonce_odr unnamed_addr addrspace(1) constant {{.*}}, comdat, align 8{{$}}
+
+// D has a key function that is defined in this translation unit so its vtable is
+// defined in the translation unit.
+// CHECK-DAG: @_ZTV1D ={{.*}} unnamed_addr addrspace(1) constant
+// CHECK-DAG: @_ZTS1D ={{.*}} addrspace(1) constant
+// CHECK-DAG: @_ZTI1D ={{.*}} addrspace(1) constant
+
+// E<char> is an explicit specialization with a key function defined
+// in this translation unit, so its vtable should have external
+// linkage.
+// CHECK-DAG: @_ZTV1EIcE ={{.*}} unnamed_addr addrspace(1) constant
+// CHECK-DAG: @_ZTS1EIcE ={{.*}} addrspace(1) constant
+// CHECK-DAG: @_ZTI1EIcE ={{.*}} addrspace(1) constant
+
+// E<short> is an explicit template instantiation with a key function
+// defined in this translation unit, so its vtable should have
+// weak_odr linkage.
+// CHECK-DAG: @_ZTV1EIsE = weak_odr unnamed_addr addrspace(1) constant {{.*}}, comdat,
+// CHECK-DAG: @_ZTS1EIsE = weak_odr addrspace(1) constant {{.*}}, comdat, align 1{{$}}
+// CHECK-DAG: @_ZTI1EIsE = weak_odr addrspace(1) constant {{.*}}, comdat, align 8{{$}}
+
+// F<short> is an explicit template instantiation without a key
+// function, so its vtable should have weak_odr linkage
+// CHECK-DAG: @_ZTV1FIsE = weak_odr unnamed_addr addrspace(1) constant {{.*}}, comdat,
+// CHECK-DAG: @_ZTS1FIsE = weak_odr addrspace(1) constant {{.*}}, comdat, align 1{{$}}
+// CHECK-DAG: @_ZTI1FIsE = weak_odr addrspace(1) constant {{.*}}, comdat, align 8{{$}}
+
+// E<long> is an implicit template instantiation with a key function
+// defined in this translation unit, so its vtable should have
+// linkonce_odr linkage.
+// CHECK-DAG: @_ZTV1EIlE = linkonce_odr unnamed_addr addrspace(1) constant {{.*}}, comdat,
+// CHECK-DAG: @_ZTS1EIlE = linkonce_odr addrspace(1) constant {{.*}}, comdat, align 1{{$}}
+// CHECK-DAG: @_ZTI1EIlE = linkonce_odr addrspace(1) constant {{.*}}, comdat, align 8{{$}}
+
+// F<long> is an implicit template instantiation with no key function,
+// so its vtable should have linkonce_odr linkage.
+// CHECK-DAG: @_ZTV1FIlE = linkonce_odr unnamed_addr addrspace(1) constant {{.*}}, comdat,
+// CHECK-DAG: @_ZTS1FIlE = linkonce_odr addrspace(1) constant {{.*}}, comdat, align 1{{$}}
+// CHECK-DAG: @_ZTI1FIlE = linkonce_odr addrspace(1) constant {{.*}}, comdat, align 8{{$}}
+
+// F<int> is an explicit template instantiation declaration without a
+// key function, so its vtable should have external linkage.
+// CHECK-DAG: @_ZTV1FIiE = external unnamed_addr addrspace(1) constant
+// CHECK-OPT-DAG: @_ZTV1FIiE = available_externally unnamed_addr addrspace(1) constant
+
+// E<int> is an explicit template instantiation declaration. It has a
+// key function is not instantiated, so we know that vtable definition
+// will be generated in TU where key function will be defined
+// so we can mark it as external (without optimizations) and
+// available_externally (with optimizations) because all of the inline
+// virtual functions have been emitted.
+// CHECK-DAG: @_ZTV1EIiE = external unnamed_addr addrspace(1) constant
+// CHECK-OPT-DAG: @_ZTV1EIiE = available_externally unnamed_addr addrspace(1) constant
+
+// The anonymous struct for e has no linkage, so the vtable should have
+// internal linkage.
+// CHECK-DAG: @"_ZTV3$_0" = internal unnamed_addr addrspace(1) constant
+// CHECK-DAG: @"_ZTS3$_0" = internal addrspace(1) constant
+// CHECK-DAG: @"_ZTI3$_0" = internal addrspace(1) constant
+
+// The A vtable should have internal linkage since it is inside an anonymous
+// namespace.
+// CHECK-DAG: @_ZTVN12_GLOBAL__N_11AE = internal unnamed_addr addrspace(1) constant
+// CHECK-DAG: @_ZTSN12_GLOBAL__N_11AE = internal addrspace(1) constant
+// CHECK-DAG: @_ZTIN12_GLOBAL__N_11AE = internal addrspace(1) constant
+
+// F<char> is an explicit specialization without a key function, so
+// its vtable should have linkonce_odr linkage.
+// CHECK-DAG: @_ZTV1FIcE = linkonce_odr unnamed_addr addrspace(1) constant {{.*}}, comdat,
+// CHECK-DAG: @_ZTS1FIcE = linkonce_odr addrspace(1) constant {{.*}}, comdat, align 1{{$}}
+// CHECK-DAG: @_ZTI1FIcE = linkonce_odr addrspace(1) constant {{.*}}, comdat, align 8{{$}}
+
+// CHECK-DAG: @_ZTV1GIiE = linkonce_odr unnamed_addr addrspace(1) constant {{.*}}, comdat,
+template <typename T>
+class G {
+public:
+  G() {}
+  virtual void f0();
+  virtual void f1();
+};
+template <>
+void G<int>::f1() {}
+template <typename T>
+void G<T>::f0() {}
+void G_f0()  { new G<int>(); }
+
+// H<int> has a key function without a body but it's a template instantiation
+// so its VTable must be emitted.
+// CHECK-DAG: @_ZTV1HIiE = linkonce_odr unnamed_addr addrspace(1) constant {{.*}}, comdat,
+template <typename T>
+class H {
+public:
+  virtual ~H();
+};
+
+void use_H() {
+  H<int> h;
+}
+
+// I<int> has an explicit instantiation declaration and needs a VTT and
+// construction vtables.
+
+// CHECK-DAG: @_ZTV1IIiE = external unnamed_addr addrspace(1) constant
+// CHECK-DAG: @_ZTT1IIiE = external unnamed_addr addrspace(1) constant
+// CHECK-NOT: @_ZTC1IIiE
+//
+// CHECK-OPT-DAG: @_ZTV1IIiE = available_externally unnamed_addr addrspace(1) constant
+// CHECK-OPT-DAG: @_ZTT1IIiE = available_externally unnamed_addr addrspace(1) constant
+struct VBase1 { virtual void f(); }; struct VBase2 : virtual VBase1 {};
+template<typename T>
+struct I : VBase2 {};
+extern template struct I<int>;
+I<int> i;

clang/test/CodeGenCXX/vtable-pointer-initialization-address-space.cpp

+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -std=c++11 -emit-llvm -o - | FileCheck %s
+
+struct Field {
+  Field();
+  ~Field();
+};
+
+struct Base {
+  Base();
+  ~Base();
+};
+
+struct A : Base {
+  A();
+  ~A();
+
+  virtual void f();
+
+  Field field;
+};
+
+// CHECK-LABEL: define{{.*}} void @_ZN1AC2Ev(ptr {{[^,]*}} %this) unnamed_addr
+// CHECK: call void @_ZN4BaseC2Ev(
+// CHECK: store ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTV1A, i32 0, inrange i32 0, i32 2)
+// CHECK: call void @_ZN5FieldC1Ev(
+// CHECK: ret void
+A::A() { }
+
+// CHECK-LABEL: define{{.*}} void @_ZN1AD2Ev(ptr {{[^,]*}} %this) unnamed_addr
+// CHECK: store ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTV1A, i32 0, inrange i32 0, i32 2)
+// CHECK: call void @_ZN5FieldD1Ev(
+// CHECK: call void @_ZN4BaseD2Ev(
+// CHECK: ret void
+A::~A() { }
+
+struct B : Base {
+  virtual void f();
+
+  Field field;
+};
+
+void f() { B b; }
+
+// CHECK-LABEL: define linkonce_odr void @_ZN1BC1Ev(ptr {{[^,]*}} %this) unnamed_addr
+// CHECK: call void @_ZN1BC2Ev(
+
+// CHECK-LABEL: define linkonce_odr void @_ZN1BD1Ev(ptr {{[^,]*}} %this) unnamed_addr
+// CHECK: call void @_ZN1BD2Ev(
+
+// CHECK-LABEL: define linkonce_odr void @_ZN1BC2Ev(ptr {{[^,]*}} %this) unnamed_addr
+// CHECK: call void @_ZN4BaseC2Ev(
+// CHECK: store ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTV1B, i32 0, inrange i32 0, i32 2)
+// CHECK: call void @_ZN5FieldC1Ev
+// CHECK: ret void
+
+// CHECK-LABEL: define linkonce_odr void @_ZN1BD2Ev(ptr {{[^,]*}} %this) unnamed_addr
+// CHECK: store ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTV1B, i32 0, inrange i32 0, i32 2)
+// CHECK: call void @_ZN5FieldD1Ev(
+// CHECK: call void @_ZN4BaseD2Ev(
+// CHECK: ret void

clang/test/CodeGenCXX/vtt-address-space.cpp

 // RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -std=c++11 -emit-llvm -o - | FileCheck %s
-// This is temporarily disabled as it requires fixing typeinfo & vptr handling
-// as well; it will be enabled once those fixes are in.
-// XFAIL: *
 
 // This is the sample from the C++ Itanium ABI, p2.6.2.
 namespace Test {
   class A1 { int i; };
   class A2 { int i; virtual void f(); };
   class V1 : public A1, public A2 { int i; };
   class B1 { int i; };
   class B2 { int i; };
   class V2 : public B1, public B2, public virtual V1 { int i; };
   class V3 { virtual void g(); };
   class C1 : public virtual V1 { int i; };
   class C2 : public virtual V3, virtual V2 { int i; };
   class X1 { int i; };
   class C3 : public X1 { int i; };
   class D : public C1, public C2, public C3 { int i;  };
 
   D d;
 }
 
-// CHECK: @_ZTTN4Test1DE = linkonce_odr unnamed_addr addrspace(1) constant [13 x ptr] [ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [5 x ptr], [7 x ptr], [4 x ptr], [3 x ptr] }, ptr addrspace(1) @_ZTVN4Test1DE, i32 0, inrange i32 0, i32 5) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [3 x ptr], [4 x ptr] }, ptr addrspace(1) @_ZTCN4Test1DE0_NS_2C1E, i32 0, inrange i32 0, i32 3) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [3 x ptr], [4 x ptr] }, ptr addrspace(1) @_ZTCN4Test1DE0_NS_2C1E, i32 0, inrange i32 1, i32 3) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [7 x ptr], [3 x ptr], [4 x ptr] }, ptr addrspace(1) @_ZTCN4Test1DE16_NS_2C2E, i32 0, inrange i32 0, i32 6) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [7 x ptr], [3 x ptr], [4 x ptr] }, ptr addrspace(1) @_ZTCN4Test1DE16_NS_2C2E, i32 0, inrange i32 0, i32 6) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [7 x ptr], [3 x ptr], [4 x ptr] }, ptr addrspace(1) @_ZTCN4Test1DE16_NS_2C2E, i32 0, inrange i32 1, i32 3) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [7 x ptr], [3 x ptr], [4 x ptr] }, ptr addrspace(1) @_ZTCN4Test1DE16_NS_2C2E, i32 0, inrange i32 2, i32 3) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [5 x ptr], [7 x ptr], [4 x ptr], [3 x ptr] }, ptr addrspace(1) @_ZTVN4Test1DE, i32 0, inrange i32 2, i32 3) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [5 x ptr], [7 x ptr], [4 x ptr], [3 x ptr] }, ptr addrspace(1) @_ZTVN4Test1DE, i32 0, inrange i32 1, i32 6) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [5 x ptr], [7 x ptr], [4 x ptr], [3 x ptr] }, ptr addrspace(1) @_ZTVN4Test1DE, i32 0, inrange i32 1, i32 6) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [5 x ptr], [7 x ptr], [4 x ptr], [3 x ptr] }, ptr addrspace(1) @_ZTVN4Test1DE, i32 0, inrange i32 3, i32 3) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [3 x ptr], [4 x ptr] }, ptr addrspace(1) @_ZTCN4Test1DE64_NS_2V2E, i32 0, inrange i32 0, i32 3) to ptr), ptr addrspacecast (ptr addrspace(1) getelementptr inbounds ({ [3 x ptr], [4 x ptr] }, ptr addrspace(1) @_ZTCN4Test1DE64_NS_2V2E, i32 0, inrange i32 1, i32 3) to ptr)], comdat, align 8
-// CHECK: call void @_ZN4Test2V2C2Ev(ptr noundef nonnull align 8 dereferenceable(20) %2, ptr addrspace(1) noundef getelementptr inbounds ([13 x ptr], ptr addrspace(1) @_ZTTN4Test1DE, i64 0, i64 11))
-// CHECK: call void @_ZN4Test2C1C2Ev(ptr noundef nonnull align 8 dereferenceable(12) %this1, ptr addrspace(1) noundef getelementptr inbounds ([13 x ptr], ptr addrspace(1) @_ZTTN4Test1DE, i64 0, i64 1))
-// CHECK: call void @_ZN4Test2C2C2Ev(ptr noundef nonnull align 8 dereferenceable(12) %3, ptr addrspace(1) noundef getelementptr inbounds ([13 x ptr], ptr addrspace(1) @_ZTTN4Test1DE, i64 0, i64 3))
-// CHECK-NEXT: define linkonce_odr void @_ZN4Test2V2C2Ev(ptr noundef nonnull align 8 dereferenceable(20) %this, ptr addrspace(1) noundef %vtt)
-// CHECK-NEXT: define linkonce_odr void @_ZN4Test2C1C2Ev(ptr noundef nonnull align 8 dereferenceable(12) %this, ptr addrspace(1) noundef %vtt)
-// CHECK-NEXT: define linkonce_odr void @_ZN4Test2C2C2Ev(ptr noundef nonnull align 8 dereferenceable(12) %this, ptr addrspace(1) noundef %vtt)
+// CHECK: linkonce_odr unnamed_addr addrspace(1) constant [13 x ptr addrspace(1)] [ptr addrspace(1) getelementptr inbounds ({ [5 x ptr addrspace(1)], [7 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN4Test1DE, i32 0, inrange i32 0, i32 5), ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN4Test1DE0_NS_2C1E, i32 0, inrange i32 0, i32 3), ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN4Test1DE0_NS_2C1E, i32 0, inrange i32 1, i32 3), ptr addrspace(1) getelementptr inbounds ({ [7 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN4Test1DE16_NS_2C2E, i32 0, inrange i32 0, i32 6), ptr addrspace(1) getelementptr inbounds ({ [7 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN4Test1DE16_NS_2C2E, i32 0, inrange i32 0, i32 6), ptr addrspace(1) getelementptr inbounds ({ [7 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN4Test1DE16_NS_2C2E, i32 0, inrange i32 1, i32 3), ptr addrspace(1) getelementptr inbounds ({ [7 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN4Test1DE16_NS_2C2E, i32 0, inrange i32 2, i32 3), ptr addrspace(1) getelementptr inbounds ({ [5 x ptr addrspace(1)], [7 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN4Test1DE, i32 0, inrange i32 2, i32 3), ptr addrspace(1) getelementptr inbounds ({ [5 x ptr addrspace(1)], [7 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN4Test1DE, i32 0, inrange i32 1, i32 6), ptr addrspace(1) getelementptr inbounds ({ [5 x ptr addrspace(1)], [7 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN4Test1DE, i32 0, inrange i32 1, i32 6), ptr addrspace(1) getelementptr inbounds ({ [5 x ptr addrspace(1)], [7 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN4Test1DE, i32 0, inrange i32 3, i32 3), ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN4Test1DE64_NS_2V2E, i32 0, inrange i32 0, i32 3), ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN4Test1DE64_NS_2V2E, i32 0, inrange i32 1, i32 3)], comdat, align 8
+// CHECK: call void @_ZN4Test2V2C2Ev(ptr noundef nonnull align 8 dereferenceable(20) %2, ptr addrspace(1) noundef getelementptr inbounds ([13 x ptr addrspace(1)], ptr addrspace(1) @_ZTTN4Test1DE, i64 0, i64 11))
+// CHECK: call void @_ZN4Test2C1C2Ev(ptr noundef nonnull align 8 dereferenceable(12) %this1, ptr addrspace(1) noundef getelementptr inbounds ([13 x ptr addrspace(1)], ptr addrspace(1) @_ZTTN4Test1DE, i64 0, i64 1))
+// CHECK: call void @_ZN4Test2C2C2Ev(ptr noundef nonnull align 8 dereferenceable(12) %3, ptr addrspace(1) noundef getelementptr inbounds ([13 x ptr addrspace(1)], ptr addrspace(1) @_ZTTN4Test1DE, i64 0, i64 3))
+// CHECK: define linkonce_odr void @_ZN4Test2V2C2Ev(ptr noundef nonnull align 8 dereferenceable(20) %this, ptr addrspace(1) noundef %vtt)
+// CHECK: define linkonce_odr void @_ZN4Test2C1C2Ev(ptr noundef nonnull align 8 dereferenceable(12) %this, ptr addrspace(1) noundef %vtt)
+// CHECK: define linkonce_odr void @_ZN4Test2C2C2Ev(ptr noundef nonnull align 8 dereferenceable(12) %this, ptr addrspace(1) noundef %vtt)

clang/test/CodeGenCXX/vtt-layout-address-space.cpp

+// RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -std=c++11 -emit-llvm -o - | FileCheck %s
+
+// Test1::B should just have a single entry in its VTT, which points to the vtable.
+namespace Test1 {
+struct A { };
+
+struct B : virtual A {
+  virtual void f();
+};
+
+void B::f() { }
+}
+
+// Check that we don't add a secondary virtual pointer for Test2::A, since Test2::A doesn't have any virtual member functions or bases.
+namespace Test2 {
+  struct A { };
+
+  struct B : A { virtual void f(); };
+  struct C : virtual B { };
+
+  C c;
+}
+
+// This is the sample from the C++ Itanium ABI, p2.6.2.
+namespace Test3 {
+  class A1 { int i; };
+  class A2 { int i; virtual void f(); };
+  class V1 : public A1, public A2 { int i; };
+  class B1 { int i; };
+  class B2 { int i; };
+  class V2 : public B1, public B2, public virtual V1 { int i; };
+  class V3 {virtual void g(); };
+  class C1 : public virtual V1 { int i; };
+  class C2 : public virtual V3, virtual V2 { int i; };
+  class X1 { int i; };
+  class C3 : public X1 { int i; };
+  class D : public C1, public C2, public C3 { int i;  };
+
+  D d;
+}
+
+// This is the sample from the C++ Itanium ABI, p2.6.2, with the change suggested
+// (making A2 a virtual base of V1)
+namespace Test4 {
+  class A1 { int i; };
+  class A2 { int i; virtual void f(); };
+  class V1 : public A1, public virtual A2 { int i; };
+  class B1 { int i; };
+  class B2 { int i; };
+  class V2 : public B1, public B2, public virtual V1 { int i; };
+  class V3 {virtual void g(); };
+  class C1 : public virtual V1 { int i; };
+  class C2 : public virtual V3, virtual V2 { int i; };
+  class X1 { int i; };
+  class C3 : public X1 { int i; };
+  class D : public C1, public C2, public C3 { int i;  };
+
+  D d;
+}
+
+namespace Test5 {
+  struct A {
+    virtual void f() = 0;
+    virtual void anchor();
+  };
+
+  void A::anchor() {
+  }
+}
+
+namespace Test6 {
+  struct A {
+    virtual void f() = delete;
+    virtual void anchor();
+  };
+
+  void A::anchor() {
+  }
+}
+
+// CHECK: @_ZTTN5Test11BE ={{.*}} unnamed_addr addrspace(1) constant [1 x ptr addrspace(1)] [ptr addrspace(1) getelementptr inbounds ({ [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test11BE, i32 0, inrange i32 0, i32 3)]
+// CHECK: @_ZTVN5Test51AE ={{.*}} unnamed_addr addrspace(1) constant { [4 x ptr addrspace(1)] } { [4 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) @_ZTIN5Test51AE, ptr addrspace(1) addrspacecast (ptr @__cxa_pure_virtual to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @_ZN5Test51A6anchorEv to ptr addrspace(1))] }
+// CHECK: @_ZTVN5Test61AE ={{.*}} unnamed_addr addrspace(1) constant { [4 x ptr addrspace(1)] } { [4 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) @_ZTIN5Test61AE, ptr addrspace(1) addrspacecast (ptr @__cxa_deleted_virtual to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @_ZN5Test61A6anchorEv to ptr addrspace(1))] }
+// CHECK: @_ZTTN5Test21CE = linkonce_odr unnamed_addr addrspace(1) constant [2 x ptr addrspace(1)] [ptr addrspace(1) getelementptr inbounds ({ [5 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test21CE, i32 0, inrange i32 0, i32 4), ptr addrspace(1) getelementptr inbounds ({ [5 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test21CE, i32 0, inrange i32 0, i32 4)]
+// CHECK: @_ZTTN5Test31DE = linkonce_odr unnamed_addr addrspace(1) constant [13 x ptr addrspace(1)] [ptr addrspace(1) getelementptr inbounds ({ [5 x ptr addrspace(1)], [7 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test31DE, i32 0, inrange i32 0, i32 5), ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test31DE0_NS_2C1E, i32 0, inrange i32 0, i32 3), ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test31DE0_NS_2C1E, i32 0, inrange i32 1, i32 3), ptr addrspace(1) getelementptr inbounds ({ [7 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test31DE16_NS_2C2E, i32 0, inrange i32 0, i32 6), ptr addrspace(1) getelementptr inbounds ({ [7 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test31DE16_NS_2C2E, i32 0, inrange i32 0, i32 6), ptr addrspace(1) getelementptr inbounds ({ [7 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test31DE16_NS_2C2E, i32 0, inrange i32 1, i32 3), ptr addrspace(1) getelementptr inbounds ({ [7 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test31DE16_NS_2C2E, i32 0, inrange i32 2, i32 3), ptr addrspace(1) getelementptr inbounds ({ [5 x ptr addrspace(1)], [7 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test31DE, i32 0, inrange i32 2, i32 3), ptr addrspace(1) getelementptr inbounds ({ [5 x ptr addrspace(1)], [7 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test31DE, i32 0, inrange i32 1, i32 6), ptr addrspace(1) getelementptr inbounds ({ [5 x ptr addrspace(1)], [7 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test31DE, i32 0, inrange i32 1, i32 6), ptr addrspace(1) getelementptr inbounds ({ [5 x ptr addrspace(1)], [7 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test31DE, i32 0, inrange i32 3, i32 3), ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test31DE64_NS_2V2E, i32 0, inrange i32 0, i32 3), ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test31DE64_NS_2V2E, i32 0, inrange i32 1, i32 3)]
+// CHECK: @_ZTVN5Test41DE = linkonce_odr unnamed_addr addrspace(1) constant { [6 x ptr addrspace(1)], [8 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)], [4 x ptr addrspace(1)] } { [6 x ptr addrspace(1)] [ptr addrspace(1) inttoptr (i64 72 to ptr addrspace(1)), ptr addrspace(1) inttoptr (i64 16 to ptr addrspace(1)), ptr addrspace(1) inttoptr (i64 56 to ptr addrspace(1)), ptr addrspace(1) inttoptr (i64 40 to ptr addrspace(1)), ptr addrspace(1) null, ptr addrspace(1) @_ZTIN5Test41DE], [8 x ptr addrspace(1)] [ptr addrspace(1) inttoptr (i64 40 to ptr addrspace(1)), ptr addrspace(1) inttoptr (i64 24 to ptr addrspace(1)), ptr addrspace(1) inttoptr (i64 56 to ptr addrspace(1)), ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) inttoptr (i64 -16 to ptr addrspace(1)), ptr addrspace(1) @_ZTIN5Test41DE, ptr addrspace(1) addrspacecast (ptr @_ZN5Test42V31gEv to ptr addrspace(1))], [3 x ptr addrspace(1)] [ptr addrspace(1) inttoptr (i64 16 to ptr addrspace(1)), ptr addrspace(1) inttoptr (i64 -40 to ptr addrspace(1)), ptr addrspace(1) @_ZTIN5Test41DE], [4 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) inttoptr (i64 -56 to ptr addrspace(1)), ptr addrspace(1) @_ZTIN5Test41DE, ptr addrspace(1) addrspacecast (ptr @_ZN5Test42A21fEv to ptr addrspace(1))], [4 x ptr addrspace(1)] [ptr addrspace(1) inttoptr (i64 -16 to ptr addrspace(1)), ptr addrspace(1) inttoptr (i64 -32 to ptr addrspace(1)), ptr addrspace(1) inttoptr (i64 -72 to ptr addrspace(1)), ptr addrspace(1) @_ZTIN5Test41DE] }
+// CHECK: @_ZTTN5Test41DE = linkonce_odr unnamed_addr addrspace(1) constant [19 x ptr addrspace(1)] [ptr addrspace(1) getelementptr inbounds ({ [6 x ptr addrspace(1)], [8 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test41DE, i32 0, inrange i32 0, i32 6), ptr addrspace(1) getelementptr inbounds ({ [4 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE0_NS_2C1E, i32 0, inrange i32 0, i32 4), ptr addrspace(1) getelementptr inbounds ({ [4 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE0_NS_2C1E, i32 0, inrange i32 1, i32 3), ptr addrspace(1) getelementptr inbounds ({ [4 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE0_NS_2C1E, i32 0, inrange i32 2, i32 3), ptr addrspace(1) getelementptr inbounds ({ [8 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE16_NS_2C2E, i32 0, inrange i32 0, i32 7), ptr addrspace(1) getelementptr inbounds ({ [8 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE16_NS_2C2E, i32 0, inrange i32 0, i32 7), ptr addrspace(1) getelementptr inbounds ({ [8 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE16_NS_2C2E, i32 0, inrange i32 1, i32 4), ptr addrspace(1) getelementptr inbounds ({ [8 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE16_NS_2C2E, i32 0, inrange i32 2, i32 3), ptr addrspace(1) getelementptr inbounds ({ [8 x ptr addrspace(1)], [4 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE16_NS_2C2E, i32 0, inrange i32 3, i32 3), ptr addrspace(1) getelementptr inbounds ({ [6 x ptr addrspace(1)], [8 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test41DE, i32 0, inrange i32 2, i32 3), ptr addrspace(1) getelementptr inbounds ({ [6 x ptr addrspace(1)], [8 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test41DE, i32 0, inrange i32 3, i32 3), ptr addrspace(1) getelementptr inbounds ({ [6 x ptr addrspace(1)], [8 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test41DE, i32 0, inrange i32 1, i32 7), ptr addrspace(1) getelementptr inbounds ({ [6 x ptr addrspace(1)], [8 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test41DE, i32 0, inrange i32 1, i32 7), ptr addrspace(1) getelementptr inbounds ({ [6 x ptr addrspace(1)], [8 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTVN5Test41DE, i32 0, inrange i32 4, i32 4), ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE40_NS_2V1E, i32 0, inrange i32 0, i32 3), ptr addrspace(1) getelementptr inbounds ({ [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE40_NS_2V1E, i32 0, inrange i32 1, i32 3), ptr addrspace(1) getelementptr inbounds ({ [4 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE72_NS_2V2E, i32 0, inrange i32 0, i32 4), ptr addrspace(1) getelementptr inbounds ({ [4 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE72_NS_2V2E, i32 0, inrange i32 1, i32 3), ptr addrspace(1) getelementptr inbounds ({ [4 x ptr addrspace(1)], [3 x ptr addrspace(1)], [4 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTCN5Test41DE72_NS_2V2E, i32 0, inrange i32 2, i32 3)]
+// CHECK: declare void @__cxa_pure_virtual() unnamed_addr
+// CHECK: declare void @__cxa_deleted_virtual() unnamed_addr

clang/test/Headers/hip-header.hip

 // RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
 // RUN:   -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
 // RUN:   -DHIP_VERSION_MAJOR=4 -DHIP_VERSION_MINOR=5 \
 // RUN:   -fsanitize=address -disable-llvm-passes -D__HIPCC_RTC__ \
 // RUN:   | FileCheck -check-prefixes=MALLOC-ASAN %s
 
 // expected-no-diagnostics
 
+// Check handling of overriden, implicitly __host__ dtor (should emit as a
+// nullptr to global)
+
+struct vbase {
+    virtual ~vbase();
+};
+
+template<typename T>
+struct vderived : public vbase {
+    ~vderived();
+};
+
+template struct vderived<void>;
+
+// CHECK: @_ZTV8vderivedIvE = weak_odr unnamed_addr addrspace(1) constant { [4 x ptr addrspace(1)] } zeroinitializer, comdat, align 8
+
 // Check support for pure and deleted virtual functions
 struct base {
   __host__
   __device__
   virtual void pv() = 0;
   __host__
   __device__
   virtual void dv() = delete;
 };
 struct derived:base {
   __host__
   __device__
   virtual void pv() override {};
 };
 __device__ void test_vf() {
     derived d;
 }
-// CHECK: @_ZTV7derived = linkonce_odr unnamed_addr addrspace(1) constant { [4 x ptr] } { [4 x ptr] [ptr null, ptr null, ptr @_ZN7derived2pvEv, ptr @__cxa_deleted_virtual] }, comdat, align 8
-// CHECK: @_ZTV4base = linkonce_odr unnamed_addr addrspace(1) constant { [4 x ptr] } { [4 x ptr] [ptr null, ptr null, ptr @__cxa_pure_virtual, ptr @__cxa_deleted_virtual] }, comdat, align 8
-
+// CHECK: @_ZTV7derived = linkonce_odr unnamed_addr addrspace(1) constant { [4 x ptr addrspace(1)] } { [4 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr @_ZN7derived2pvEv to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @__cxa_deleted_virtual to ptr addrspace(1))] }, comdat, align 8
+// CHECK: @_ZTV4base = linkonce_odr unnamed_addr addrspace(1) constant { [4 x ptr addrspace(1)] } { [4 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr @__cxa_pure_virtual to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @__cxa_deleted_virtual to ptr addrspace(1))] }, comdat, align 8
 // CHECK: define{{.*}}void @__cxa_pure_virtual()
 // CHECK: define{{.*}}void @__cxa_deleted_virtual()
 
 struct Number {
   __device__ Number(float _x) : x(_x) {}
   float x;
 };