[clang] Avoid duplicating ProgramAddressSpace in TargetInfo. NFCI

[arichardson]

Summary

This value was added to clang/Basic in D111566, but is only used during
codegen, where we can use the LLVM IR DataLayout instead. I noticed this
because the downstream CHERI targets would have to also set this value
for AArch64/RISC-V/MIPS. Instead of duplicating more information between
LLVM IR and Clang, this patch moves getTargetAddressSpace(QualType T) to
CodeGenTypes, where we can consult the DataLayout.

[arichardson]

clang-format

[eandrews]

Functionally this looks ok to me. However I am not sure if CodeGenTypes is the 'right' place for this function to live, considering that other functions with similar functionality are in ASTContext - including overloads of getTargetAddressSpace( ). @erichkeane @aaron.ballman could you please take a look?

[arichardson]

In D138296#3937224, @eandrews wrote:

Functionally this looks ok to me. However I am not sure if CodeGenTypes is the 'right' place for this function to live, considering that other functions with similar functionality are in ASTContext - including overloads of getTargetAddressSpace( ). @erichkeane @aaron.ballman could you please take a look?

My view is that the parts that interact with LLVM IR should really live in CodeGen/ and not Basic/ or AST/. I will see how difficult it would be to move the remaining target (LLVM IR) address space handling code to CodeGen/

[rjmccall]

In D138296#3937486, @arichardson wrote:

In D138296#3937224, @eandrews wrote:

Functionally this looks ok to me. However I am not sure if CodeGenTypes is the 'right' place for this function to live, considering that other functions with similar functionality are in ASTContext - including overloads of getTargetAddressSpace( ). @erichkeane @aaron.ballman could you please take a look?

My view is that the parts that interact with LLVM IR should really live in CodeGen/ and not Basic/ or AST/. I will see how difficult it would be to move the remaining target (LLVM IR) address space handling code to CodeGen/

Yeah, I don't think there's a good reason for some of the address-space stuff that currently lives in Basic to be there instead of in CodeGen. Basic/AST need to understand what address spaces exist, their sizes, and what relationships they have with each other, and that's it.

[arichardson]

In D138296#3937599, @rjmccall wrote:

In D138296#3937486, @arichardson wrote:

In D138296#3937224, @eandrews wrote:

Functionally this looks ok to me. However I am not sure if CodeGenTypes is the 'right' place for this function to live, considering that other functions with similar functionality are in ASTContext - including overloads of getTargetAddressSpace( ). @erichkeane @aaron.ballman could you please take a look?

My view is that the parts that interact with LLVM IR should really live in CodeGen/ and not Basic/ or AST/. I will see how difficult it would be to move the remaining target (LLVM IR) address space handling code to CodeGen/

Yeah, I don't think there's a good reason for some of the address-space stuff that currently lives in Basic to be there instead of in CodeGen. Basic/AST need to understand what address spaces exist, their sizes, and what relationships they have with each other, and that's it.

Thanks for looking at this - does this mean you are happy for me to commit this change?

[rjmccall]

I think this is fine. Most of the patch is changing calls to getTargetAddressSpace to be internal to IRGen, which, as mentioned, I think is the right move.

I do think that if we're going to support multiple program address spaces (which seems to be a goal) that we'll eventually want an AST-level concept of the default program address space, but perpetuating the use of target ASes at the AST level isn't the right way to approach that, so this is still the right first step.

clang/include/clang/AST/ASTContext.h

public:
   /// \p LHS < \p RHS, return -1.
   int getFloatingTypeOrder(QualType LHS, QualType RHS) const;
 
   /// Compare the rank of two floating point types as above, but compare equal
   /// if both types have the same floating-point semantics on the target (i.e.
   /// long double and double on AArch64 will return 0).
   int getFloatingTypeSemanticOrder(QualType LHS, QualType RHS) const;
 
-  unsigned getTargetAddressSpace(QualType T) const;
-
   unsigned getTargetAddressSpace(LangAS AS) const;
 
   LangAS getLangASForBuiltinAddressSpace(unsigned AS) const;
 
   /// Get target-dependent integer value for null pointer which is used for
   /// constant folding.
   uint64_t getTargetNullPointerValue(QualType QT) const;
 

clang/include/clang/Basic/TargetInfo.h

protected:
   unsigned char MaxAtomicPromoteWidth, MaxAtomicInlineWidth;
   unsigned short SimdDefaultAlign;
   std::string DataLayoutString;
   const char *UserLabelPrefix;
   const char *MCountName;
   unsigned char RegParmMax, SSERegParmMax;
   TargetCXXABI TheCXXABI;
   const LangASMap *AddrSpaceMap;
-  unsigned ProgramAddrSpace;
 
   mutable StringRef PlatformName;
   mutable VersionTuple PlatformMinVersion;
 
   unsigned HasAlignMac68kSupport : 1;
   unsigned RealTypeUsesObjCFPRetMask : llvm::BitWidth<FloatModeKind>;
   unsigned ComplexLongDoubleUsesFP2Ret : 1;
 
virtual unsigned getExnObjectAlignment() const {
     return getDefaultAlignForAttributeAligned();
   }
 
   /// Return the size of intmax_t and uintmax_t for this target, in bits.
   unsigned getIntMaxTWidth() const {
     return getTypeWidth(IntMaxType);
   }
 
-  /// Return the address space for functions for the given target.
-  unsigned getProgramAddressSpace() const { return ProgramAddrSpace; }
-
   // Return the size of unwind_word for this target.
   virtual unsigned getUnwindWordWidth() const {
     return getPointerWidth(LangAS::Default);
   }
 
   /// Return the "preferred" register width on this target.
   virtual unsigned getRegisterWidth() const {
     // Currently we assume the register width on the target matches the pointer

clang/lib/AST/ASTContext.cpp

uint64_t ASTContext::getTargetNullPointerValue(QualType QT) const {
   if (QT->getUnqualifiedDesugaredType()->isNullPtrType())
     AS = LangAS::Default;
   else
     AS = QT->getPointeeType().getAddressSpace();
 
   return getTargetInfo().getNullPointerValue(AS);
 }
 
-unsigned ASTContext::getTargetAddressSpace(QualType T) const {
-  // Return the address space for the type. If the type is a
-  // function type without an address space qualifier, the
-  // program address space is used. Otherwise, the target picks
-  // the best address space based on the type information
-  return T->isFunctionType() && !T.hasAddressSpace()
-             ? getTargetInfo().getProgramAddressSpace()
-             : getTargetAddressSpace(T.getAddressSpace());
-}
-
 unsigned ASTContext::getTargetAddressSpace(LangAS AS) const {
   if (isTargetAddressSpace(AS))
     return toTargetAddressSpace(AS);
   else
     return (*AddrSpaceMap)[(unsigned)AS];
 }
 
 bool ASTContext::hasSameExpr(const Expr *X, const Expr *Y) const {

clang/lib/Basic/TargetInfo.cpp

TargetInfo::TargetInfo(const llvm::Triple &T) : Triple(T) {
 
   // Default to an unknown platform name.
   PlatformName = "unknown";
   PlatformMinVersion = VersionTuple();
 
   MaxOpenCLWorkGroupSize = 1024;
 
   MaxBitIntWidth.reset();
-
-  ProgramAddrSpace = 0;
 }
 
 // Out of line virtual dtor for TargetInfo.
 TargetInfo::~TargetInfo() {}
 
 void TargetInfo::resetDataLayout(StringRef DL, const char *ULP) {
   DataLayoutString = DL.str();
   UserLabelPrefix = ULP;

clang/lib/Basic/Targets/AVR.h

AVRTargetInfo(const llvm::Triple &Triple, const TargetOptions &)
     SizeType = UnsignedInt;
     PtrDiffType = SignedInt;
     IntPtrType = SignedInt;
     Char16Type = UnsignedInt;
     WIntType = SignedInt;
     Int16Type = SignedInt;
     Char32Type = UnsignedLong;
     SigAtomicType = SignedChar;
-    ProgramAddrSpace = 1;
     resetDataLayout("e-P1-p:16:8-i8:8-i16:8-i32:8-i64:8-f32:8-f64:8-n8-a:8");
   }
 
   void getTargetDefines(const LangOptions &Opts,
                         MacroBuilder &Builder) const override;
 
   ArrayRef<Builtin::Info> getTargetBuiltins() const override { return None; }
 

clang/lib/CodeGen/CGCall.cpp

case ABIArgInfo::Direct:
     resultType = retAI.getCoerceToType();
     break;
 
   case ABIArgInfo::InAlloca:
     if (retAI.getInAllocaSRet()) {
       // sret things on win32 aren't void, they return the sret pointer.
       QualType ret = FI.getReturnType();
       llvm::Type *ty = ConvertType(ret);
-      unsigned addressSpace = Context.getTargetAddressSpace(ret);
+      unsigned addressSpace = CGM.getTypes().getTargetAddressSpace(ret);
       resultType = llvm::PointerType::get(ty, addressSpace);
     } else {
       resultType = llvm::Type::getVoidTy(getLLVMContext());
     }
     break;
 
   case ABIArgInfo::Indirect:
   case ABIArgInfo::Ignore:
     resultType = llvm::Type::getVoidTy(getLLVMContext());
     break;
 
   case ABIArgInfo::CoerceAndExpand:
     resultType = retAI.getUnpaddedCoerceAndExpandType();
     break;
   }
 
   ClangToLLVMArgMapping IRFunctionArgs(getContext(), FI, true);
   SmallVector<llvm::Type*, 8> ArgTypes(IRFunctionArgs.totalIRArgs());
 
   // Add type for sret argument.
   if (IRFunctionArgs.hasSRetArg()) {
     QualType Ret = FI.getReturnType();
     llvm::Type *Ty = ConvertType(Ret);
-    unsigned AddressSpace = Context.getTargetAddressSpace(Ret);
+    unsigned AddressSpace = CGM.getTypes().getTargetAddressSpace(Ret);
     ArgTypes[IRFunctionArgs.getSRetArgNo()] =
         llvm::PointerType::get(Ty, AddressSpace);
   }
 
   // Add type for inalloca argument.
   if (IRFunctionArgs.hasInallocaArg()) {
     auto ArgStruct = FI.getArgStruct();
     assert(ArgStruct);
void CodeGenModule::ConstructAttributeList(StringRef Name,
 
   if (!IsThunk) {
     // FIXME: fix this properly, https://reviews.llvm.org/D100388
     if (const auto *RefTy = RetTy->getAs<ReferenceType>()) {
       QualType PTy = RefTy->getPointeeType();
       if (!PTy->isIncompleteType() && PTy->isConstantSizeType())
         RetAttrs.addDereferenceableAttr(
             getMinimumObjectSize(PTy).getQuantity());
-      if (getContext().getTargetAddressSpace(PTy) == 0 &&
+      if (getTypes().getTargetAddressSpace(PTy) == 0 &&
           !CodeGenOpts.NullPointerIsValid)
         RetAttrs.addAttribute(llvm::Attribute::NonNull);
       if (PTy->isObjectType()) {
         llvm::Align Alignment =
             getNaturalPointeeTypeAlignment(RetTy).getAsAlign();
         RetAttrs.addAlignmentAttr(Alignment);
       }
     }
if (FI.isInstanceMethod() && !IRFunctionArgs.hasInallocaArg() &&
     assert(IRArgs.second == 1 && "Expected only a single `this` pointer.");
 
     llvm::AttrBuilder Attrs(getLLVMContext());
 
     QualType ThisTy =
         FI.arg_begin()->type.castAs<PointerType>()->getPointeeType();
 
     if (!CodeGenOpts.NullPointerIsValid &&
-        getContext().getTargetAddressSpace(FI.arg_begin()->type) == 0) {
+        getTypes().getTargetAddressSpace(FI.arg_begin()->type) == 0) {
       Attrs.addAttribute(llvm::Attribute::NonNull);
       Attrs.addDereferenceableAttr(getMinimumObjectSize(ThisTy).getQuantity());
     } else {
       // FIXME dereferenceable should be correct here, regardless of
       // NullPointerIsValid. However, dereferenceable currently does not always
       // respect NullPointerIsValid and may imply nonnull and break the program.
       // See https://reviews.llvm.org/D66618 for discussions.
       Attrs.addDereferenceableOrNullAttr(
case ABIArgInfo::InAlloca:
       continue;
     }
 
     if (const auto *RefTy = ParamType->getAs<ReferenceType>()) {
       QualType PTy = RefTy->getPointeeType();
       if (!PTy->isIncompleteType() && PTy->isConstantSizeType())
         Attrs.addDereferenceableAttr(
             getMinimumObjectSize(PTy).getQuantity());
-      if (getContext().getTargetAddressSpace(PTy) == 0 &&
+      if (getTypes().getTargetAddressSpace(PTy) == 0 &&
           !CodeGenOpts.NullPointerIsValid)
         Attrs.addAttribute(llvm::Attribute::NonNull);
       if (PTy->isObjectType()) {
         llvm::Align Alignment =
             getNaturalPointeeTypeAlignment(ParamType).getAsAlign();
         Attrs.addAlignmentAttr(Alignment);
       }
     }
case ABIArgInfo::Direct: {
             // For C99 VLAs with the static keyword, we don't know the size so
             // we can't use the dereferenceable attribute, but in addrspace(0)
             // we know that it must be nonnull.
             if (ArrTy->getSizeModifier() == VariableArrayType::Static) {
               QualType ETy = ArrTy->getElementType();
               llvm::Align Alignment =
                   CGM.getNaturalTypeAlignment(ETy).getAsAlign();
               AI->addAttrs(llvm::AttrBuilder(getLLVMContext()).addAlignmentAttr(Alignment));
-              if (!getContext().getTargetAddressSpace(ETy) &&
+              if (!getTypes().getTargetAddressSpace(ETy) &&
                   !CGM.getCodeGenOpts().NullPointerIsValid)
                 AI->addAttr(llvm::Attribute::NonNull);
             }
           }
 
           // Set `align` attribute if any.
           const auto *AVAttr = PVD->getAttr<AlignValueAttr>();
           if (!AVAttr)

clang/lib/CodeGen/CGDebugInfo.cpp

llvm::DIType *CGDebugInfo::CreatePointerLikeType(llvm::dwarf::Tag Tag,
                                                  const Type *Ty,
                                                  QualType PointeeTy,
                                                  llvm::DIFile *Unit) {
   // Bit size, align and offset of the type.
   // Size is always the size of a pointer.
   uint64_t Size = CGM.getContext().getTypeSize(Ty);
   auto Align = getTypeAlignIfRequired(Ty, CGM.getContext());
   Optional<unsigned> DWARFAddressSpace = CGM.getTarget().getDWARFAddressSpace(
-      CGM.getContext().getTargetAddressSpace(PointeeTy));
+      CGM.getTypes().getTargetAddressSpace(PointeeTy));
 
   SmallVector<llvm::Metadata *, 4> Annots;
   auto *BTFAttrTy = dyn_cast<BTFTagAttributedType>(PointeeTy);
   while (BTFAttrTy) {
     StringRef Tag = BTFAttrTy->getAttr()->getBTFTypeTag();
     if (!Tag.empty()) {
       llvm::Metadata *Ops[2] = {
           llvm::MDString::get(CGM.getLLVMContext(), StringRef("btf_type_tag")),
llvm::DILocalVariable *CGDebugInfo::EmitDeclare(const VarDecl *VD,
   }
   SmallVector<uint64_t, 13> Expr;
   llvm::DINode::DIFlags Flags = llvm::DINode::FlagZero;
   if (VD->isImplicit())
     Flags |= llvm::DINode::FlagArtificial;
 
   auto Align = getDeclAlignIfRequired(VD, CGM.getContext());
 
-  unsigned AddressSpace = CGM.getContext().getTargetAddressSpace(VD->getType());
+  unsigned AddressSpace = CGM.getTypes().getTargetAddressSpace(VD->getType());
   AppendAddressSpaceXDeref(AddressSpace, Expr);
 
   // If this is implicit parameter of CXXThis or ObjCSelf kind, then give it an
   // object pointer flag.
   if (const auto *IPD = dyn_cast<ImplicitParamDecl>(VD)) {
     if (IPD->getParameterKind() == ImplicitParamDecl::CXXThis ||
         IPD->getParameterKind() == ImplicitParamDecl::ObjCSelf)
       Flags |= llvm::DINode::FlagObjectPointer;
llvm::DILocalVariable *CGDebugInfo::EmitDeclare(const BindingDecl *BD,
   llvm::DIType *Ty = getOrCreateType(BD->getType(), Unit);
 
   // If there is no debug info for this type then do not emit debug info
   // for this variable.
   if (!Ty)
     return nullptr;
 
   auto Align = getDeclAlignIfRequired(BD, CGM.getContext());
-  unsigned AddressSpace = CGM.getContext().getTargetAddressSpace(BD->getType());
+  unsigned AddressSpace = CGM.getTypes().getTargetAddressSpace(BD->getType());
 
   SmallVector<uint64_t, 3> Expr;
   AppendAddressSpaceXDeref(AddressSpace, Expr);
 
   // Clang stores the sret pointer provided by the caller in a static alloca.
   // Use DW_OP_deref to tell the debugger to load the pointer and treat it as
   // the address of the variable.
   if (UsePointerValue) {
if (T->isUnionType() && DeclName.empty()) {
     const RecordDecl *RD = T->castAs<RecordType>()->getDecl();
     assert(RD->isAnonymousStructOrUnion() &&
            "unnamed non-anonymous struct or union?");
     GVE = CollectAnonRecordDecls(RD, Unit, LineNo, LinkageName, Var, DContext);
   } else {
     auto Align = getDeclAlignIfRequired(D, CGM.getContext());
 
     SmallVector<uint64_t, 4> Expr;
-    unsigned AddressSpace =
-        CGM.getContext().getTargetAddressSpace(D->getType());
+    unsigned AddressSpace = CGM.getTypes().getTargetAddressSpace(D->getType());
     if (CGM.getLangOpts().CUDA && CGM.getLangOpts().CUDAIsDevice) {
       if (D->hasAttr<CUDASharedAttr>())
         AddressSpace =
             CGM.getContext().getTargetAddressSpace(LangAS::cuda_shared);
       else if (D->hasAttr<CUDAConstantAttr>())
         AddressSpace =
             CGM.getContext().getTargetAddressSpace(LangAS::cuda_constant);
     }

clang/lib/CodeGen/CGDeclCXX.cpp

void CodeGenFunction::EmitCXXGlobalVarDeclInit(const VarDecl &D,
   // __device__ void foo() {
   //   __shared__ StructWithCtor s;
   //   ...
   // }
   //
   // For example, in the above CUDA code, the static local variable s has a
   // "shared" address space qualifier, but the constructor of StructWithCtor
   // expects "this" in the "generic" address space.
-  unsigned ExpectedAddrSpace = getContext().getTargetAddressSpace(T);
+  unsigned ExpectedAddrSpace = getTypes().getTargetAddressSpace(T);
   unsigned ActualAddrSpace = GV->getAddressSpace();
   llvm::Constant *DeclPtr = GV;
   if (ActualAddrSpace != ExpectedAddrSpace) {
     llvm::PointerType *PTy = llvm::PointerType::getWithSamePointeeType(
         GV->getType(), ExpectedAddrSpace);
     DeclPtr = llvm::ConstantExpr::getAddrSpaceCast(DeclPtr, PTy);
   }
 

clang/lib/CodeGen/CGExpr.cpp

if (E->isNonOdrUse() == NOUR_Constant &&
 
       Address Addr = Address::invalid();
       if (!VD->getType()->isReferenceType()) {
         // Spill the constant value to a global.
         Addr = CGM.createUnnamedGlobalFrom(*VD, Val,
                                            getContext().getDeclAlign(VD));
         llvm::Type *VarTy = getTypes().ConvertTypeForMem(VD->getType());
         auto *PTy = llvm::PointerType::get(
-            VarTy, getContext().getTargetAddressSpace(VD->getType()));
+            VarTy, getTypes().getTargetAddressSpace(VD->getType()));
         Addr = Builder.CreatePointerBitCastOrAddrSpaceCast(Addr, PTy, VarTy);
       } else {
         // Should we be using the alignment of the constant pointer we emitted?
         CharUnits Alignment =
             CGM.getNaturalTypeAlignment(E->getType(),
                                         /* BaseInfo= */ nullptr,
                                         /* TBAAInfo= */ nullptr,
                                         /* forPointeeType= */ true);

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

assert(NativeParam != TargetParam &&
          NativeParam->getType()->isReferenceType() &&
          "Native arg must not be the same as target arg.");
   Address LocalAddr = CGF.GetAddrOfLocalVar(TargetParam);
   QualType NativeParamType = NativeParam->getType();
   QualifierCollector QC;
   const Type *NonQualTy = QC.strip(NativeParamType);
   QualType NativePointeeTy = cast<ReferenceType>(NonQualTy)->getPointeeType();
   unsigned NativePointeeAddrSpace =
-      CGF.getContext().getTargetAddressSpace(NativePointeeTy);
+      CGF.getTypes().getTargetAddressSpace(NativePointeeTy);
   QualType TargetTy = TargetParam->getType();
   llvm::Value *TargetAddr = CGF.EmitLoadOfScalar(
       LocalAddr, /*Volatile=*/false, TargetTy, SourceLocation());
   // First cast to generic.
   TargetAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
       TargetAddr, llvm::PointerType::getWithSamePointeeType(
           cast<llvm::PointerType>(TargetAddr->getType()), /*AddrSpace=*/0));
   // Cast from generic to native address space.

clang/lib/CodeGen/CodeGenModule.cpp

ConstantAddress CodeGenModule::GetWeakRefReference(const ValueDecl *VD) {
   assert(AA && "No alias?");
 
   CharUnits Alignment = getContext().getDeclAlign(VD);
   llvm::Type *DeclTy = getTypes().ConvertTypeForMem(VD->getType());
 
   // See if there is already something with the target's name in the module.
   llvm::GlobalValue *Entry = GetGlobalValue(AA->getAliasee());
   if (Entry) {
-    unsigned AS = getContext().getTargetAddressSpace(VD->getType());
+    unsigned AS = getTypes().getTargetAddressSpace(VD->getType());
     auto Ptr = llvm::ConstantExpr::getBitCast(Entry, DeclTy->getPointerTo(AS));
     return ConstantAddress(Ptr, DeclTy, Alignment);
   }
 
   llvm::Constant *Aliasee;
   if (isa<llvm::FunctionType>(DeclTy))
     Aliasee = GetOrCreateLLVMFunction(AA->getAliasee(), DeclTy,
                                       GlobalDecl(cast<FunctionDecl>(VD)),
void CodeGenModule::emitCPUDispatchDefinition(GlobalDecl GD) {
   StringRef ResolverName = getMangledName(GD);
   UpdateMultiVersionNames(GD, FD, ResolverName);
 
   llvm::Type *ResolverType;
   GlobalDecl ResolverGD;
   if (getTarget().supportsIFunc()) {
     ResolverType = llvm::FunctionType::get(
         llvm::PointerType::get(DeclTy,
-                               Context.getTargetAddressSpace(FD->getType())),
+                               getTypes().getTargetAddressSpace(FD->getType())),
         false);
   }
   else {
     ResolverType = DeclTy;
     ResolverGD = GD;
   }
 
   auto *ResolverFunc = cast<llvm::Function>(GetOrCreateLLVMFunction(
llvm::Constant *CodeGenModule::GetOrCreateMultiVersionResolver(GlobalDecl GD) {
   // creation until the end of the TU.
   if (FD->isTargetMultiVersion() || FD->isTargetClonesMultiVersion())
     MultiVersionFuncs.push_back(GD);
 
   // For cpu_specific, don't create an ifunc yet because we don't know if the
   // cpu_dispatch will be emitted in this translation unit.
   if (getTarget().supportsIFunc() && !FD->isCPUSpecificMultiVersion()) {
     llvm::Type *ResolverType = llvm::FunctionType::get(
-        llvm::PointerType::get(
-            DeclTy, getContext().getTargetAddressSpace(FD->getType())),
+        llvm::PointerType::get(DeclTy,
+                               getTypes().getTargetAddressSpace(FD->getType())),
         false);
     llvm::Constant *Resolver = GetOrCreateLLVMFunction(
         MangledName + ".resolver", ResolverType, GlobalDecl{},
         /*ForVTable=*/false);
     llvm::GlobalIFunc *GIF =
         llvm::GlobalIFunc::create(DeclTy, 0, getMultiversionLinkage(*this, GD),
                                   "", Resolver, &getModule());
     GIF->setName(ResolverName);

clang/lib/CodeGen/CodeGenTypes.h

public:  // These are internal details of CGT that shouldn't be used externally.
 
   bool isRecordLayoutComplete(const Type *Ty) const;
   bool noRecordsBeingLaidOut() const {
     return RecordsBeingLaidOut.empty();
   }
   bool isRecordBeingLaidOut(const Type *Ty) const {
     return RecordsBeingLaidOut.count(Ty);
   }
-
+  unsigned getTargetAddressSpace(QualType T) const;
 };
 
 }  // end namespace CodeGen
 }  // end namespace clang
 
 #endif

clang/lib/CodeGen/CodeGenTypes.cpp

case Type::Complex: {
     ResultType = llvm::StructType::get(EltTy, EltTy);
     break;
   }
   case Type::LValueReference:
   case Type::RValueReference: {
     const ReferenceType *RTy = cast<ReferenceType>(Ty);
     QualType ETy = RTy->getPointeeType();
     llvm::Type *PointeeType = ConvertTypeForMem(ETy);
-    unsigned AS = Context.getTargetAddressSpace(ETy);
+    unsigned AS = getTargetAddressSpace(ETy);
     ResultType = llvm::PointerType::get(PointeeType, AS);
     break;
   }
   case Type::Pointer: {
     const PointerType *PTy = cast<PointerType>(Ty);
     QualType ETy = PTy->getPointeeType();
     llvm::Type *PointeeType = ConvertTypeForMem(ETy);
     if (PointeeType->isVoidTy())
       PointeeType = llvm::Type::getInt8Ty(getLLVMContext());
-    unsigned AS = Context.getTargetAddressSpace(ETy);
+    unsigned AS = getTargetAddressSpace(ETy);
     ResultType = llvm::PointerType::get(PointeeType, AS);
     break;
   }
 
   case Type::VariableArray: {
     const VariableArrayType *A = cast<VariableArrayType>(Ty);
     assert(A->getIndexTypeCVRQualifiers() == 0 &&
            "FIXME: We only handle trivial array types so far!");
bool CodeGenTypes::isZeroInitializable(QualType T) {
 
   // Everything else is okay.
   return true;
 }
 
 bool CodeGenTypes::isZeroInitializable(const RecordDecl *RD) {
   return getCGRecordLayout(RD).isZeroInitializable();
 }
+
+unsigned CodeGenTypes::getTargetAddressSpace(QualType T) const {
+  // Return the address space for the type. If the type is a
+  // function type without an address space qualifier, the
+  // program address space is used. Otherwise, the target picks
+  // the best address space based on the type information
+  return T->isFunctionType() && !T.hasAddressSpace()
+             ? getDataLayout().getProgramAddressSpace()
+             : getContext().getTargetAddressSpace(T.getAddressSpace());
+}