[WIP] Supporting C++ based kernel languages on AMDGPU Target

[yaxunl]

Summary

In C++ based languages targeting GPUGPU, e.g. HCC, OpenMP or Cuda, it is desirable to keep the AST free of explicit address space qualifier for simplified parsing and type checking. We can assume that all the pointers are generic pointers which can pointing to any real address spaces in the target IR. This results in a consistent and clean AST and therefore IR since every pointer is a generic pointer.

In Clang codegen, when we need to emit global variables or allocas in real address space, we cast the resulting pointer to a generic pointer and put them in the address map. Later on, when we need to refer to the address, we always get a generic pointer. In this way, we can keep the translation of AST consistent.

Currently Clang assumes that the address space for a generic pointer (a pointer which can pointing to any real memory region) is 0, which is correct for x86 and nvptx target, but is not correct for some other targets, e.g., amdgcn target, for which the generic address space is 4.

This patch introduce a virtual function TargetInfo::getDefaultTargetAddressSpace which allows each target to specify the default target address space based on language options. It also introduces ASTContext::getTargetGlobalAddressSpace and ASTContext::getTargetConstantAddressSpace as an abstraction for a unified approach to support different languages on address space aware targets, e.g. amdgcn and nvptx.

It also fixes various places in codegen for incorrect assumptions about default address space.

With this change, Clang will be able to generate correct IR for C++ based kernel languages e.g. HCC, OpenMP and Cuda for a greater range of targets.

It will not change the emitted IR for targets using 0 as generic address space. For these targets, since alloca returns generic pointer, they usually require some LLVM passes to cast alloca to real address space. This patch can help those targets generate correct alloca in Clang codegen.

This patch was co-authored by Yaxun (Sam) Liu and Wen-Heng (Jack) Chung.

[rjmccall]

What address space should local variables be in?

[yaxunl]

In D27627#618597, @rjmccall wrote:

What address space should local variables be in?

should always be 0 since that's llvm convention.

[rjmccall]

Ok, I'm not understanding this, then.

I declare a local variable:

int x;

According to you, x should be in the generic address space, i.e. the private address space. And at the LLVM IR level, this will be implemented with an AllocaInst, which always creates memory in the private address space, i.e. LLVM address space 0.

But Sema believes that the l-value 'x' has type int, without an address space, and therefore that '&x' has type int*, again without an explicit address space. And you want to say that that type actually points into a different address space, and to make sure that that pointer is in LLVM address space 4.

Something about this does not work.

[rjmccall]

I suspect that what you actually want to do in order to implement HCC is change SemaType so that 'int*' is implicitly interpreted as 'int __something *'. But I don't know how not having explicit address spaces actually works in the HCC language design, given that it's based on C++.

[rjmccall]

Because, notably, if you do that, then an attempt to pass &x as an int* will fail, which means this isn't really C++ anymore... and yet that appears to be exactly what you want.

[yaxunl]

Because, notably, if you do that, then an attempt to pass &x as an int* will fail, which means this isn't really C++ anymore... and yet that appears to be exactly what you want.

How about when generating alloca instruction, I insert addrspacecast to the default address space, then everything is in default address space.

[rjmccall]

In D27627#619928, @yaxunl wrote:

Because, notably, if you do that, then an attempt to pass &x as an int* will fail, which means this isn't really C++ anymore... and yet that appears to be exactly what you want.

How about when generating alloca instruction, I insert addrspacecast to the default address space, then everything is in default address space.

My question is really about your language design. You have multiple address spaces in the implementation, but you're (apparently?) pretending that they're all part of the same address space in the source language. How is that expected to work? Does the default address space embed all other address spaces?

[yaxunl]

In D27627#621473, @rjmccall wrote:

In D27627#619928, @yaxunl wrote:

Because, notably, if you do that, then an attempt to pass &x as an int* will fail, which means this isn't really C++ anymore... and yet that appears to be exactly what you want.

How about when generating alloca instruction, I insert addrspacecast to the default address space, then everything is in default address space.

My question is really about your language design. You have multiple address spaces in the implementation, but you're (apparently?) pretending that they're all part of the same address space in the source language. How is that expected to work? Does the default address space embed all other address spaces?

Yes the default address space embeds all other address spaces.

[yaxunl]

Cast alloca to default address space.

[yaxunl]

In D27627#621533, @yaxunl wrote:

In D27627#621473, @rjmccall wrote:

In D27627#619928, @yaxunl wrote:

Because, notably, if you do that, then an attempt to pass &x as an int* will fail, which means this isn't really C++ anymore... and yet that appears to be exactly what you want.

How about when generating alloca instruction, I insert addrspacecast to the default address space, then everything is in default address space.

My question is really about your language design. You have multiple address spaces in the implementation, but you're (apparently?) pretending that they're all part of the same address space in the source language. How is that expected to work? Does the default address space embed all other address spaces?

Yes the default address space embeds all other address spaces.

Since the default address space (generic address space 4) can represent all other address spaces, we want to use it uniformly. For any pointer which is not in the default address space, e.g. alloca, we want to cast it to default address space and then use it. Therefore, in the address space agnostic languages, we can assume everything is in the default address space. Our backend is able to handle load/store in the default address space.

[yaxunl]

Synch up to trunk.
Change bitcast to pointer cast in varous places to accommodation addr space difference.

[yaxunl]

Fix invalid bitcast for global variables.
Add run line for amdgcn target for a Cuda lit test.

[yaxunl]

Fixed invalid bitcasts due to vtable.
Added run line for amdgcn target to more Cuda lit tests.

[kzhuravl]

Once you rebase, can you also update AMDGPUTargetInfo::getDWARFAddressSpace to use AddrSpaceKind?

[yaxunl]

We consider to switch to use 0 as generic address space and abandon this since 0 as generic address space is more agreeable to both LLVM and Clang's convention.

[rjmccall]

Okay. That does seem more sensible than trying to change everything around the other way.

include/clang/AST/ASTContext.h

   /// \brief Return a real floating point or a complex type (based on
   /// \p typeDomain/\p typeSize).
   ///
   /// \param typeDomain a real floating point or complex type.
   /// \param typeSize a real floating point or complex type.
   QualType getFloatingTypeOfSizeWithinDomain(QualType typeSize,
                                              QualType typeDomain) const;
 
-  unsigned getTargetAddressSpace(QualType T) const {
-    return getTargetAddressSpace(T.getQualifiers());
-  }
+  unsigned getTargetAddressSpace(QualType T) const;
 
-  unsigned getTargetAddressSpace(Qualifiers Q) const {
-    return getTargetAddressSpace(Q.getAddressSpace());
-  }
+  unsigned getTargetAddressSpace(Qualifiers Q) const;
 
   unsigned getTargetAddressSpace(unsigned AS) const {
     if (AS < LangAS::Offset || AS >= LangAS::Offset + LangAS::Count)
       return AS;
     else
       return (*AddrSpaceMap)[AS - LangAS::Offset];
   }
 
   /// Get target-dependent integer value for null pointer which is used for
   /// constant folding.
   uint64_t getTargetNullPointerValue(QualType QT) const;
 
+  unsigned getTargetDefaultAddressSpace() const;
+
+  /// The target address space corresponding to OpenCL constant address space
+  /// CUDA constant specifier.
+  unsigned getTargetConstantAddressSpace() const;
+
+  /// The target address space corresponding to OpenCL global address space
+  /// or CUDA device specifier.
+  unsigned getTargetGlobalAddressSpace() const;
+
   bool addressSpaceMapManglingFor(unsigned AS) const {
     return AddrSpaceMapMangling || 
            AS < LangAS::Offset || 
            AS >= LangAS::Offset + LangAS::Count;
   }
 
 private:
   // Helper for integer ordering

include/clang/Basic/TargetInfo.h

public:
   }
 
   /// \brief Return the maximum width of pointers on this target.
   virtual uint64_t getMaxPointerWidth() const {
     return PointerWidth;
   }
 
   /// \brief Get integer value for null pointer.
-  /// \param AddrSpace address space of pointee in source language.
+  /// \param AddrSpace target address space of pointee.
   virtual uint64_t getNullPointerValue(unsigned AddrSpace) const {
     return 0;
   }
 
+  /// The target address space corresponding to OpenCL constant address space
+  /// CUDA constant specifier.
+  virtual unsigned getConstantAddressSpace() const {
+    return 0;
+  }
+
+  /// The target address space corresponding to OpenCL global address space
+  /// or CUDA device specifier.
+  virtual unsigned getGlobalAddressSpace() const {
+    return 0;
+  }
+
   /// \brief Return the size of '_Bool' and C++ 'bool' for this target, in bits.
   unsigned getBoolWidth() const { return BoolWidth; }
 
   /// \brief Return the alignment of '_Bool' and C++ 'bool' for this target.
   unsigned getBoolAlign() const { return BoolAlign; }
 
   unsigned getCharWidth() const { return 8; } // FIXME
   unsigned getCharAlign() const { return 8; } // FIXME
public:
   virtual const char *getStaticInitSectionSpecifier() const {
     return nullptr;
   }
 
   const LangAS::Map &getAddressSpaceMap() const {
     return *AddrSpaceMap;
   }
 
+  virtual unsigned getDefaultTargetAddressSpace(const LangOptions &Opt) const {
+    return 0;
+  }
+
   /// \brief Retrieve the name of the platform as it is used in the
   /// availability attribute.
   StringRef getPlatformName() const { return PlatformName; }
 
   /// \brief Retrieve the minimum desired version of the platform, to
   /// which the program should be compiled.
   VersionTuple getPlatformMinVersion() const { return PlatformMinVersion; }
 

lib/AST/ASTContext.cpp

   }
   return (MethodDecl->isVariadic() == MethodImpl->isVariadic());
   
 }
 
 uint64_t ASTContext::getTargetNullPointerValue(QualType QT) const {
   unsigned AS;
   if (QT->getUnqualifiedDesugaredType()->isNullPtrType())
-    AS = 0;
+    AS = getTargetInfo().getDefaultTargetAddressSpace(LangOpts);
   else
-    AS = QT->getPointeeType().getAddressSpace();
+    AS = getTargetAddressSpace(QT->getPointeeType());
 
   return getTargetInfo().getNullPointerValue(AS);
 }
 
+unsigned ASTContext::getTargetDefaultAddressSpace() const {
+  return getTargetInfo().getDefaultTargetAddressSpace(LangOpts);
+}
+
+unsigned ASTContext::getTargetConstantAddressSpace() const {
+  return getTargetInfo().getConstantAddressSpace();
+}
+
+unsigned ASTContext::getTargetGlobalAddressSpace() const {
+  return getTargetInfo().getGlobalAddressSpace();
+}
+
+unsigned ASTContext::getTargetAddressSpace(QualType T) const {
+  if (T.isNull())
+    return getTargetDefaultAddressSpace();
+  if (T->isFunctionType() &&
+      !T.getQualifiers().hasAddressSpace())
+    return 0;
+  return getTargetAddressSpace(T.getQualifiers());
+}
+
+unsigned ASTContext::getTargetAddressSpace(Qualifiers Q) const {
+  return Q.hasAddressSpace()
+             ? getTargetAddressSpace(Q.getAddressSpace())
+             : getTargetDefaultAddressSpace();
+}
+
 // Explicitly instantiate this in case a Redeclarable<T> is used from a TU that
 // doesn't include ASTContext.h
 template
 clang::LazyGenerationalUpdatePtr<
     const Decl *, Decl *, &ExternalASTSource::CompleteRedeclChain>::ValueType
 clang::LazyGenerationalUpdatePtr<
     const Decl *, Decl *, &ExternalASTSource::CompleteRedeclChain>::makeValue(
         const clang::ASTContext &Ctx, Decl *Value);

lib/Basic/TargetInfo.cpp

if (Opts.OpenCL) {
 
     HalfFormat = &llvm::APFloat::IEEEhalf();
     FloatFormat = &llvm::APFloat::IEEEsingle();
     LongDoubleFormat = &llvm::APFloat::IEEEquad();
   }
 
   if (Opts.NewAlignOverride)
     NewAlign = Opts.NewAlignOverride * getCharWidth();
+
+  if (getTriple().getArch() == llvm::Triple::amdgcn) {
+    auto DefAddr = getDefaultTargetAddressSpace(Opts);
+    // AMDGPUTargetInfo only implements getPointerWidthV and assumes
+    // pointers are self-aligned.
+    PointerWidth = PointerAlign = getPointerWidthV(DefAddr);
+  }
 }
 
 bool TargetInfo::initFeatureMap(
     llvm::StringMap<bool> &Features, DiagnosticsEngine &Diags, StringRef CPU,
     const std::vector<std::string> &FeatureVec) const {
   for (const auto &F : FeatureVec) {
     StringRef Name = F;
     // Apply the feature via the target.

lib/Basic/Targets.cpp

 };
 
 const char *const NVPTXTargetInfo::GCCRegNames[] = {"r0"};
 
 ArrayRef<const char *> NVPTXTargetInfo::getGCCRegNames() const {
   return llvm::makeArrayRef(GCCRegNames);
 }
 
-static const unsigned AMDGPUAddrSpaceMap[] = {
-  1,    // opencl_global
-  3,    // opencl_local
-  2,    // opencl_constant
-  4,    // opencl_generic
-  1,    // cuda_device
-  2,    // cuda_constant
-  3     // cuda_shared
-};
-
 // If you edit the description strings, make sure you update
 // getPointerWidthV().
 
 static const char *const DataLayoutStringR600 =
   "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128"
   "-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64";
 
 static const char *const DataLayoutStringSI =
   "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32"
   "-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128"
   "-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64";
 
 class AMDGPUTargetInfo final : public TargetInfo {
+  static const unsigned AddrSpaceMap_[7];
   static const Builtin::Info BuiltinInfo[];
   static const char * const GCCRegNames[];
 
+  enum AddrSpaceKind {
+    AS_Private = 0,
+    AS_Global = 1,
+    AS_Constant = 2,
+    AS_Local = 3,
+    AS_Generic = 4
+  };
+
   /// \brief The GPU profiles supported by the AMDGPU target.
   enum GPUKind {
     GK_NONE,
     GK_R600,
     GK_R600_DOUBLE_OPS,
     GK_R700,
     GK_R700_DOUBLE_OPS,
     GK_EVERGREEN,

       hasFP64 = true;
       hasFMAF = true;
       hasLDEXPF = true;
     }
 
     resetDataLayout(getTriple().getArch() == llvm::Triple::amdgcn ?
                     DataLayoutStringSI : DataLayoutStringR600);
 
-    AddrSpaceMap = &AMDGPUAddrSpaceMap;
+    AddrSpaceMap = &AddrSpaceMap_;
     UseAddrSpaceMapMangling = true;
   }
 
   uint64_t getPointerWidthV(unsigned AddrSpace) const override {
     if (GPU <= GK_CAYMAN)
       return 32;
 
     switch(AddrSpace) {

       Opts.support("cl_khr_mipmap_image");
       Opts.support("cl_khr_subgroups");
       Opts.support("cl_khr_3d_image_writes");
       Opts.support("cl_amd_media_ops");
       Opts.support("cl_amd_media_ops2");
     }
   }
 
+  unsigned
+  getDefaultTargetAddressSpace(const LangOptions &Opts) const override {
+    // OpenCL sets address space explicitly in AST. The default case (type
+    // qualifier containing no address space) represents private address space.
+    if (Opts.OpenCL)
+      return AS_Private;
+    return AS_Generic;
+  }
+
+  unsigned getConstantAddressSpace() const override {
+    return AS_Constant;
+  }
+
+  unsigned getGlobalAddressSpace() const override {
+    return AS_Global;
+  }
+
   LangAS::ID getOpenCLImageAddrSpace() const override {
     return LangAS::opencl_constant;
   }
 
   CallingConvCheckResult checkCallingConvention(CallingConv CC) const override {
     switch (CC) {
       default:
         return CCCR_Warning;
       case CC_C:
       case CC_OpenCLKernel:
         return CCCR_OK;
     }
   }
 
-  // In amdgcn target the null pointer in global, constant, and generic
-  // address space has value 0 but in private and local address space has
-  // value ~0.
+  // In amdgcn target the null pointer in local and private address spaces has
+  // value ~0 and in other address spaces has value 0.
   uint64_t getNullPointerValue(unsigned AS) const override {
-    return AS != LangAS::opencl_local && AS != 0 ? 0 : ~0;
+    return AS != AS_Local && AS != 0 ? 0 : ~0;
   }
 };
 
+const unsigned AMDGPUTargetInfo::AddrSpaceMap_[] = {
+    AS_Global,   // opencl_global
+    AS_Local,    // opencl_local
+    AS_Constant, // opencl_constant
+    AS_Generic,  // opencl_generic
+    AS_Global,   // cuda_device
+    AS_Constant, // cuda_constant
+    AS_Local     // cuda_shared
+};
+
 const Builtin::Info AMDGPUTargetInfo::BuiltinInfo[] = {
 #define BUILTIN(ID, TYPE, ATTRS)                \
   { #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, nullptr },
 #define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE)                               \
   { #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, FEATURE },
 #include "clang/Basic/BuiltinsAMDGPU.def"
 };
 const char * const AMDGPUTargetInfo::GCCRegNames[] = {

lib/CodeGen/CGBuiltin.cpp

       return RValue::get(CS.getInstruction());
     }
     break;
   }
 
   case Builtin::BI__GetExceptionInfo: {
     if (llvm::GlobalVariable *GV =
             CGM.getCXXABI().getThrowInfo(FD->getParamDecl(0)->getType()))
-      return RValue::get(llvm::ConstantExpr::getBitCast(GV, CGM.Int8PtrTy));
+      return RValue::get(llvm::ConstantExpr::getPointerCast(GV, CGM.Int8PtrTy));
     break;
   }
 
   case Builtin::BI__fastfail: {
     return RValue::get(EmitMSVCBuiltinExpr(MSVCIntrin::__fastfail, E));
     break;
   }
 

     Arg = Builder.CreatePointerCast(Arg, GenericVoidPtrTy);
     return RValue::get(Builder.CreateCall(
         CGM.CreateRuntimeFunction(
             llvm::FunctionType::get(IntTy, GenericVoidPtrTy, false),
             "__get_kernel_preferred_work_group_multiple_impl"),
         Arg));
   }
   case Builtin::BIprintf:
-    if (getTarget().getTriple().isNVPTX())
+    if (getTarget().getTriple().isNVPTX() ||
+        (getTarget().getTriple().getArch() == Triple::amdgcn &&
+         getLangOpts().CUDA))
       return EmitNVPTXDevicePrintfCallExpr(E, ReturnValue);
     break;
   case Builtin::BI__builtin_canonicalize:
   case Builtin::BI__builtin_canonicalizef:
   case Builtin::BI__builtin_canonicalizel:
     return RValue::get(emitUnaryBuiltin(*this, E, Intrinsic::canonicalize));
 
   case Builtin::BI__builtin_thread_pointer: {

lib/CodeGen/CGCall.cpp

 
   // If we're using inalloca, insert the allocation after the stack save.
   // FIXME: Do this earlier rather than hacking it in here!
   Address ArgMemory = Address::invalid();
   const llvm::StructLayout *ArgMemoryLayout = nullptr;
   if (llvm::StructType *ArgStruct = CallInfo.getArgStruct()) {
     ArgMemoryLayout = CGM.getDataLayout().getStructLayout(ArgStruct);
     llvm::Instruction *IP = CallArgs.getStackBase();
-    llvm::AllocaInst *AI;
+    llvm::Instruction *CastedAI;
     if (IP) {
       IP = IP->getNextNode();
-      AI = new llvm::AllocaInst(ArgStruct, "argmem", IP);
+      CastedAI = CreateAlloca(ArgStruct, "argmem", IP);
     } else {
-      AI = CreateTempAlloca(ArgStruct, "argmem");
+      CastedAI = CreateTempAlloca(ArgStruct, "argmem");
     }
     auto Align = CallInfo.getArgStructAlignment();
+    auto *AI = getAddrSpaceCastedAlloca(CastedAI);
     AI->setAlignment(Align.getQuantity());
     AI->setUsedWithInAlloca(true);
     assert(AI->isUsedWithInAlloca() && !AI->isStaticAlloca());
-    ArgMemory = Address(AI, Align);
+    ArgMemory = Address(CastedAI, Align);
   }
 
   // Helper function to drill into the inalloca allocation.
   auto createInAllocaStructGEP = [&](unsigned FieldIndex) -> Address {
     auto FieldOffset =
       CharUnits::fromQuantity(ArgMemoryLayout->getElementOffset(FieldIndex));
     return Builder.CreateStructGEP(ArgMemory, FieldIndex, FieldOffset);
   };

lib/CodeGen/CGClass.cpp

 
   if (!NonVirtualOffset.isZero() || VirtualOffset)
     VTableField = ApplyNonVirtualAndVirtualOffset(
         *this, VTableField, NonVirtualOffset, VirtualOffset, Vptr.VTableClass,
         Vptr.NearestVBase);
 
   // Finally, store the address point. Use the same LLVM types as the field to
   // support optimization.
+  auto DefAddr = CGM.getTarget().getDefaultTargetAddressSpace(
+      CGM.getLangOpts());
   llvm::Type *VTablePtrTy =
       llvm::FunctionType::get(CGM.Int32Ty, /*isVarArg=*/true)
-          ->getPointerTo()
-          ->getPointerTo();
-  VTableField = Builder.CreateBitCast(VTableField, VTablePtrTy->getPointerTo());
-  VTableAddressPoint = Builder.CreateBitCast(VTableAddressPoint, VTablePtrTy);
+          ->getPointerTo(DefAddr)
+          ->getPointerTo(DefAddr);
+  VTableField = Builder.CreatePointerBitCastOrAddrSpaceCast(VTableField,
+      VTablePtrTy->getPointerTo(DefAddr));
+  VTableAddressPoint = Builder.CreatePointerBitCastOrAddrSpaceCast(
+      VTableAddressPoint, VTablePtrTy);
 
   llvm::StoreInst *Store = Builder.CreateStore(VTableAddressPoint, VTableField);
   CGM.DecorateInstructionWithTBAA(Store, CGM.getTBAAInfoForVTablePtr());
   if (CGM.getCodeGenOpts().OptimizationLevel > 0 &&
       CGM.getCodeGenOpts().StrictVTablePointers)
     CGM.DecorateInstructionWithInvariantGroup(Store, Vptr.VTableClass);
 }
 

lib/CodeGen/CGDecl.cpp

     std::tie(elementCount, elementType) = getVLASize(Ty);
 
     llvm::Type *llvmTy = ConvertTypeForMem(elementType);
 
     // Allocate memory for the array.
     llvm::AllocaInst *vla = Builder.CreateAlloca(llvmTy, elementCount, "vla");
     vla->setAlignment(alignment.getQuantity());
 
-    address = Address(vla, alignment);
+    llvm::Value *V = vla;
+    auto DefaultAddr = getTarget().getDefaultTargetAddressSpace(getLangOpts());
+    if (DefaultAddr != 0) {
+      auto *DestTy =
+          llvm::PointerType::get(vla->getType()->getElementType(), DefaultAddr);
+      V = Builder.CreateAddrSpaceCast(vla, DestTy);
+    }
+
+    address = Address(V, alignment);
   }
 
   setAddrOfLocalVar(&D, address);
   emission.Addr = address;
 
   // Emit debug info for local var declaration.
   if (HaveInsertPoint())
     if (CGDebugInfo *DI = getDebugInfo()) {

       Loc = Builder.CreateBitCast(Loc, constant->getType()->getPointerTo());
       emitStoresForInitAfterMemset(constant, Loc.getPointer(),
                                    isVolatile, Builder);
     }
   } else {
     // Otherwise, create a temporary global with the initializer then
     // memcpy from the global to the alloca.
     std::string Name = getStaticDeclName(CGM, D);
-    unsigned AS = 0;
+    unsigned AS = CGM.getContext().getTargetConstantAddressSpace();
     if (getLangOpts().OpenCL) {
       AS = CGM.getContext().getTargetAddressSpace(LangAS::opencl_constant);
       BP = llvm::PointerType::getInt8PtrTy(getLLVMContext(), AS);
     }
     llvm::GlobalVariable *GV =
       new llvm::GlobalVariable(CGM.getModule(), constant->getType(), true,
                                llvm::GlobalValue::PrivateLinkage,
                                constant, Name, nullptr,

lib/CodeGen/CGDeclCXX.cpp

static void EmitDeclDestroy(CodeGenFunction &CGF, const VarDecl &D,
   // generated elsewhere which uses atexit instead, and it takes the destructor
   // directly.
   bool UsingExternalHelper = !CGM.getCodeGenOpts().CXAAtExit;
   if (Record && (CanRegisterDestructor || UsingExternalHelper)) {
     assert(!Record->hasTrivialDestructor());
     CXXDestructorDecl *dtor = Record->getDestructor();
 
     function = CGM.getAddrOfCXXStructor(dtor, StructorType::Complete);
-    argument = llvm::ConstantExpr::getBitCast(
-        addr.getPointer(), CGF.getTypes().ConvertType(type)->getPointerTo());
+    argument = llvm::ConstantExpr::getPointerCast(
+        addr.getPointer(), CGF.getTypes().getPointerTypeTo(type));
 
   // Otherwise, the standard logic requires a helper function.
   } else {
     function = CodeGenFunction(CGM)
         .generateDestroyHelper(addr, type, CGF.getDestroyer(dtorKind),
                                CGF.needsEHCleanup(dtorKind), &D);
     argument = llvm::Constant::getNullValue(CGF.Int8PtrTy);
   }
static void EmitDeclInvariant(CodeGenFunction &CGF, const VarDecl &D,
   // Overloaded address space type.
   llvm::Type *ObjectPtr[1] = {CGF.Int8PtrTy};
   llvm::Constant *InvariantStart = CGF.CGM.getIntrinsic(InvStartID, ObjectPtr);
 
   // Emit a call with the size in bytes of the object.
   CharUnits WidthChars = CGF.getContext().getTypeSizeInChars(D.getType());
   uint64_t Width = WidthChars.getQuantity();
   llvm::Value *Args[2] = { llvm::ConstantInt::getSigned(CGF.Int64Ty, Width),
-                           llvm::ConstantExpr::getBitCast(Addr, CGF.Int8PtrTy)};
+                           llvm::ConstantExpr::getPointerCast(Addr, CGF.Int8PtrTy)};
   CGF.Builder.CreateCall(InvariantStart, Args);
 }
 
 void CodeGenFunction::EmitCXXGlobalVarDeclInit(const VarDecl &D,
                                                llvm::Constant *DeclPtr,
                                                bool PerformInit) {
 
   const Expr *Init = D.getInit();

lib/CodeGen/CGException.cpp

static llvm::Constant *getPersonalityFn(CodeGenModule &CGM,
   return CGM.CreateRuntimeFunction(llvm::FunctionType::get(CGM.Int32Ty, true),
                                    Personality.PersonalityFn,
                                    llvm::AttributeSet(), /*Local=*/true);
 }
 
 static llvm::Constant *getOpaquePersonalityFn(CodeGenModule &CGM,
                                         const EHPersonality &Personality) {
   llvm::Constant *Fn = getPersonalityFn(CGM, Personality);
-  return llvm::ConstantExpr::getBitCast(Fn, CGM.Int8PtrTy);
+  return llvm::ConstantExpr::getPointerCast(Fn, CGM.Int8PtrTy);
 }
 
 /// Check whether a landingpad instruction only uses C++ features.
 static bool LandingPadHasOnlyCXXUses(llvm::LandingPadInst *LPI) {
   for (unsigned I = 0, E = LPI->getNumClauses(); I != E; ++I) {
     // Look for something that would've been returned by the ObjC
     // runtime's GetEHType() method.
     llvm::Value *Val = LPI->getClause(I)->stripPointerCasts();
if (auto *ParentAlloca = dyn_cast<llvm::AllocaInst>(ParentVar.getPointer())) {
     // localescape index.
     auto InsertPair = ParentCGF.EscapedLocals.insert(
         std::make_pair(ParentAlloca, ParentCGF.EscapedLocals.size()));
     int FrameEscapeIdx = InsertPair.first->second;
     // call i8* @llvm.localrecover(i8* bitcast(@parentFn), i8* %fp, i32 N)
     llvm::Function *FrameRecoverFn = llvm::Intrinsic::getDeclaration(
         &CGM.getModule(), llvm::Intrinsic::localrecover);
     llvm::Constant *ParentI8Fn =
-        llvm::ConstantExpr::getBitCast(ParentCGF.CurFn, Int8PtrTy);
+        llvm::ConstantExpr::getPointerCast(ParentCGF.CurFn, Int8PtrTy);
     RecoverCall = Builder.CreateCall(
         FrameRecoverFn, {ParentI8Fn, ParentFP,
                          llvm::ConstantInt::get(Int32Ty, FrameEscapeIdx)});
 
   } else {
     // If the parent didn't have an alloca, we're doing some nested outlining.
     // Just clone the existing localrecover call, but tweak the FP argument to
     // use our FP value. All other arguments are constants.
void CodeGenFunction::EmitCapturedLocals(CodeGenFunction &ParentCGF,
   llvm::Value *ParentFP = EntryFP;
   if (IsFilter) {
     // Given whatever FP the runtime provided us in EntryFP, recover the true
     // frame pointer of the parent function. We only need to do this in filters,
     // since finally funclets recover the parent FP for us.
     llvm::Function *RecoverFPIntrin =
         CGM.getIntrinsic(llvm::Intrinsic::x86_seh_recoverfp);
     llvm::Constant *ParentI8Fn =
-        llvm::ConstantExpr::getBitCast(ParentCGF.CurFn, Int8PtrTy);
+        llvm::ConstantExpr::getPointerCast(ParentCGF.CurFn, Int8PtrTy);
     ParentFP = Builder.CreateCall(RecoverFPIntrin, {ParentI8Fn, EntryFP});
   }
 
   // Create llvm.localrecover calls for all captures.
   for (const VarDecl *VD : Finder.Captures) {
     if (isa<ImplicitParamDecl>(VD)) {
       CGM.ErrorUnsupported(VD, "'this' captured by SEH");
       CXXThisValue = llvm::UndefValue::get(ConvertTypeForMem(VD->getType()));
if (CGM.getTarget().getTriple().getArch() != llvm::Triple::x86 && C &&
     return;
   }
 
   // In general, we have to emit an outlined filter function. Use the function
   // in place of the RTTI typeinfo global that C++ EH uses.
   llvm::Function *FilterFunc =
       HelperCGF.GenerateSEHFilterFunction(*this, *Except);
   llvm::Constant *OpaqueFunc =
-      llvm::ConstantExpr::getBitCast(FilterFunc, Int8PtrTy);
+      llvm::ConstantExpr::getPointerCast(FilterFunc, Int8PtrTy);
   CatchScope->setHandler(0, OpaqueFunc, createBasicBlock("__except.ret"));
 }
 
 void CodeGenFunction::ExitSEHTryStmt(const SEHTryStmt &S) {
   // Just pop the cleanup if it's a __finally block.
   if (S.getFinallyHandler()) {
     PopCleanupBlock();
     return;

lib/CodeGen/CGExpr.cpp

   if (value->getType() == destType) return value;
   return Builder.CreateBitCast(value, destType);
 }
 
 /// CreateTempAlloca - This creates a alloca and inserts it into the entry
 /// block.
 Address CodeGenFunction::CreateTempAlloca(llvm::Type *Ty, CharUnits Align,
                                           const Twine &Name) {
-  auto Alloca = CreateTempAlloca(Ty, Name);
+  auto CastedAlloca = CreateTempAlloca(Ty, Name);
+  auto *Alloca = getAddrSpaceCastedAlloca(CastedAlloca);
   Alloca->setAlignment(Align.getQuantity());
-  return Address(Alloca, Align);
+  return Address(CastedAlloca, Align);
 }
 
 /// CreateTempAlloca - This creates a alloca and inserts it into the entry
 /// block.
-llvm::AllocaInst *CodeGenFunction::CreateTempAlloca(llvm::Type *Ty,
+llvm::Instruction *CodeGenFunction::CreateTempAlloca(llvm::Type *Ty,
                                                      const Twine &Name) {
-  return new llvm::AllocaInst(Ty, nullptr, Name, AllocaInsertPt);
+  return CreateAlloca(Ty, Name, AllocaInsertPt);
+}
+
+llvm::Instruction *CodeGenFunction::CreateAlloca(llvm::Type *Ty,
+                                                 const Twine &Name,
+                                                 llvm::Instruction *InsertPos) {
+  llvm::Instruction *V = new llvm::AllocaInst(Ty, nullptr, Name, InsertPos);
+  auto DefaultAddr = getTarget().getDefaultTargetAddressSpace(getLangOpts());
+  if (DefaultAddr != 0) {
+    auto *DestTy = llvm::PointerType::get(V->getType()->getPointerElementType(),
+                                          DefaultAddr);
+    V = new llvm::AddrSpaceCastInst(V, DestTy, "", InsertPos);
+  }
+  return V;
+}
+
+llvm::AllocaInst *
+CodeGenFunction::getAddrSpaceCastedAlloca(llvm::Instruction *V) const {
+  if (auto *Cast = dyn_cast<llvm::AddrSpaceCastInst>(V))
+    return cast<llvm::AllocaInst>(Cast->getOperand(0));
+  return cast<llvm::AllocaInst>(V);
 }
 
 /// CreateDefaultAlignTempAlloca - This creates an alloca with the
 /// default alignment of the corresponding LLVM type, which is *not*
 /// guaranteed to be related in any way to the expected alignment of
 /// an AST type that might have been lowered to Ty.
 Address CodeGenFunction::CreateDefaultAlignTempAlloca(llvm::Type *Ty,
                                                       const Twine &Name) {

       assert(Adjustments.empty());
       return EmitOpaqueValueLValue(opaque);
     }
   }
 
   // Create and initialize the reference temporary.
   Address Object = createReferenceTemporary(*this, M, E);
   if (auto *Var = dyn_cast<llvm::GlobalVariable>(Object.getPointer())) {
-    Object = Address(llvm::ConstantExpr::getBitCast(
-        Var, ConvertTypeForMem(E->getType())->getPointerTo()),
+    Object = Address(llvm::ConstantExpr::getPointerCast(
+        Var, getTypes().getPointerTypeTo(E->getType())),
                      Object.getAlignment());
     // If the temporary is a global and has a constant initializer or is a
     // constant temporary that we promoted to a global, we may have already
     // initialized it.
     if (!Var->hasInitializer()) {
       Var->setInitializer(CGM.EmitNullConstant(E->getType()));
       EmitAnyExprToMem(E, Object, Qualifiers(), /*IsInit*/true);
     }

   // anything here.
   if (!E->getType()->isVariableArrayType()) {
     assert(isa<llvm::ArrayType>(Addr.getElementType()) &&
            "Expected pointer to array");
     Addr = Builder.CreateStructGEP(Addr, 0, CharUnits::Zero(), "arraydecay");
   }
 
   QualType EltType = E->getType()->castAsArrayTypeUnsafe()->getElementType();
-  return Builder.CreateElementBitCast(Addr, ConvertTypeForMem(EltType));
+  return Builder.CreatePointerBitCastOrAddrSpaceCast(Addr,
+      ConvertTypeForMem(EltType)->getPointerTo(getContext().
+          getTargetAddressSpace(E->getType())));
 }
 
 /// isSimpleArrayDecayOperand - If the specified expr is a simple decay from an
 /// array to pointer, return the array subexpression.
 static const Expr *isSimpleArrayDecayOperand(const Expr *E) {
   // If this isn't just an array->pointer decay, bail out.
   const auto *CE = dyn_cast<CastExpr>(E);
   if (!CE || CE->getCastKind() != CK_ArrayToPointerDecay)

lib/CodeGen/CGExprCXX.cpp

     CGF.EmitBlock(EndBlock);
   }
 
   return CGF.CGM.getCXXABI().EmitTypeid(CGF, SrcRecordTy, ThisPtr,
                                         StdTypeInfoPtrTy);
 }
 
 llvm::Value *CodeGenFunction::EmitCXXTypeidExpr(const CXXTypeidExpr *E) {
-  llvm::Type *StdTypeInfoPtrTy = 
-    ConvertType(E->getType())->getPointerTo();
+  llvm::Type *StdTypeInfoPtrTy = getTypes().getPointerTypeTo(E->getType());
   
   if (E->isTypeOperand()) {
     llvm::Constant *TypeInfo =
         CGM.GetAddrOfRTTIDescriptor(E->getTypeOperand(getContext()));
     return Builder.CreateBitCast(TypeInfo, StdTypeInfoPtrTy);
   }
 
   // C++ [expr.typeid]p2:

lib/CodeGen/CGExprConstant.cpp

       }
 
       C = ConstExprEmitter(*this, CGF).EmitLValue(LVBase).getPointer();
 
       // Apply offset if necessary.
       if (!Offset->isNullValue()) {
         unsigned AS = C->getType()->getPointerAddressSpace();
         llvm::Type *CharPtrTy = Int8Ty->getPointerTo(AS);
-        llvm::Constant *Casted = llvm::ConstantExpr::getBitCast(C, CharPtrTy);
+        llvm::Constant *Casted = llvm::ConstantExpr::getPointerCast(C, CharPtrTy);
         Casted = llvm::ConstantExpr::getGetElementPtr(Int8Ty, Casted, Offset);
         C = llvm::ConstantExpr::getPointerCast(Casted, C->getType());
       }
 
       // Convert to the appropriate type; this could be an lvalue for
       // an integer.
       if (isa<llvm::PointerType>(DestTy))
         return llvm::ConstantExpr::getPointerCast(C, DestTy);

lib/CodeGen/CGExprScalar.cpp

Value *EmitIntToBoolConversion(Value *V) {
     return Builder.CreateIsNotNull(V, "tobool");
   }
 
   //===--------------------------------------------------------------------===//
   //                            Visitor Methods
   //===--------------------------------------------------------------------===//
 
   Value *Visit(Expr *E) {
+    if (getenv("DBG_CG_SCALAR_EXPR")) {
+      llvm::errs() << "Expr: "; E->dump();
+    }
     ApplyDebugLocation DL(CGF, E);
-    return StmtVisitor<ScalarExprEmitter, Value*>::Visit(E);
+    auto Res = StmtVisitor<ScalarExprEmitter, Value*>::Visit(E);
+    if (getenv("DBG_CG_SCALAR_EXPR")) {
+      llvm::errs() << " => " << *Res << '\n';
+    }
+    return Res;
   }
 
   Value *VisitStmt(Stmt *S) {
     S->dump(CGF.getContext().getSourceManager());
     llvm_unreachable("Stmt can't have complex result type!");
   }
   Value *VisitExpr(Expr *S);
 

lib/CodeGen/CGGPUBuiltin.cpp

 #include "clang/Basic/Builtins.h"
 #include "llvm/IR/DataLayout.h"
 #include "llvm/IR/Instruction.h"
 #include "llvm/Support/MathExtras.h"
 
 using namespace clang;
 using namespace CodeGen;
 
-static llvm::Function *GetVprintfDeclaration(llvm::Module &M) {
-  llvm::Type *ArgTypes[] = {llvm::Type::getInt8PtrTy(M.getContext()),
-                            llvm::Type::getInt8PtrTy(M.getContext())};
+static llvm::Function *GetVprintfDeclaration(CodeGenModule &CGM) {
+  auto &M = CGM.getModule();
+  llvm::Type *ArgTypes[] = {CGM.Int8PtrTy, CGM.Int8PtrTy};
   llvm::FunctionType *VprintfFuncType = llvm::FunctionType::get(
       llvm::Type::getInt32Ty(M.getContext()), ArgTypes, false);
 
   if (auto* F = M.getFunction("vprintf")) {
     // Our CUDA system header declares vprintf with the right signature, so
     // nobody else should have been able to declare vprintf with a bogus
     // signature.
     assert(F->getFunctionType() == VprintfFuncType);

 // buf is aligned to the max of {alignof(Arg1), ...}.  Furthermore, each of the
 // args is itself aligned to its preferred alignment.
 //
 // Note that by the time this function runs, E's args have already undergone the
 // standard C vararg promotion (short -> int, float -> double, etc.).
 RValue
 CodeGenFunction::EmitNVPTXDevicePrintfCallExpr(const CallExpr *E,
                                                ReturnValueSlot ReturnValue) {
-  assert(getTarget().getTriple().isNVPTX());
+  assert(getTarget().getTriple().isNVPTX() ||
+        (getTarget().getTriple().getArch() == llvm::Triple::amdgcn &&
+         getLangOpts().CUDA));
   assert(E->getBuiltinCallee() == Builtin::BIprintf);
   assert(E->getNumArgs() >= 1); // printf always has at least one arg.
 
   const llvm::DataLayout &DL = CGM.getDataLayout();
-  llvm::LLVMContext &Ctx = CGM.getLLVMContext();
 
   CallArgList Args;
   EmitCallArgs(Args,
                E->getDirectCallee()->getType()->getAs<FunctionProtoType>(),
                E->arguments(), E->getDirectCallee(),
                /* ParamsToSkip = */ 0);
 
   // We don't know how to emit non-scalar varargs.
   if (std::any_of(Args.begin() + 1, Args.end(),
                   [](const CallArg &A) { return !A.RV.isScalar(); })) {
     CGM.ErrorUnsupported(E, "non-scalar arg to printf");
     return RValue::get(llvm::ConstantInt::get(IntTy, 0));
   }
 
   // Construct and fill the args buffer that we'll pass to vprintf.
   llvm::Value *BufferPtr;
   if (Args.size() <= 1) {
     // If there are no args, pass a null pointer to vprintf.
-    BufferPtr = llvm::ConstantPointerNull::get(llvm::Type::getInt8PtrTy(Ctx));
+    BufferPtr = llvm::ConstantPointerNull::get(CGM.Int8PtrTy);
   } else {
     llvm::SmallVector<llvm::Type *, 8> ArgTypes;
     for (unsigned I = 1, NumArgs = Args.size(); I < NumArgs; ++I)
       ArgTypes.push_back(Args[I].RV.getScalarVal()->getType());
 
     // Using llvm::StructType is correct only because printf doesn't accept
     // aggregates.  If we had to handle aggregates here, we'd have to manually
     // compute the offsets within the alloca -- we wouldn't be able to assume
     // that the alignment of the llvm type was the same as the alignment of the
     // clang type.
     llvm::Type *AllocaTy = llvm::StructType::create(ArgTypes, "printf_args");
     llvm::Value *Alloca = CreateTempAlloca(AllocaTy);
 
     for (unsigned I = 1, NumArgs = Args.size(); I < NumArgs; ++I) {
       llvm::Value *P = Builder.CreateStructGEP(AllocaTy, Alloca, I - 1);
       llvm::Value *Arg = Args[I].RV.getScalarVal();
       Builder.CreateAlignedStore(Arg, P, DL.getPrefTypeAlignment(Arg->getType()));
     }
-    BufferPtr = Builder.CreatePointerCast(Alloca, llvm::Type::getInt8PtrTy(Ctx));
+    BufferPtr = Builder.CreatePointerCast(Alloca, CGM.Int8PtrTy);
   }
 
   // Invoke vprintf and return.
-  llvm::Function* VprintfFunc = GetVprintfDeclaration(CGM.getModule());
+  llvm::Function* VprintfFunc = GetVprintfDeclaration(CGM);
   return RValue::get(
       Builder.CreateCall(VprintfFunc, {Args[0].RV.getScalarVal(), BufferPtr}));
 }

lib/CodeGen/CGOpenMPRuntime.cpp

if (!Entry) {
     if (!DefaultOpenMPPSource) {
       // Initialize default location for psource field of ident_t structure of
       // all ident_t objects. Format is ";file;function;line;column;;".
       // Taken from
       // http://llvm.org/svn/llvm-project/openmp/trunk/runtime/src/kmp_str.c
       DefaultOpenMPPSource =
           CGM.GetAddrOfConstantCString(";unknown;unknown;0;0;;").getPointer();
       DefaultOpenMPPSource =
-          llvm::ConstantExpr::getBitCast(DefaultOpenMPPSource, CGM.Int8PtrTy);
+          llvm::ConstantExpr::getPointerCast(DefaultOpenMPPSource, CGM.Int8PtrTy);
     }
 
     ConstantInitBuilder builder(CGM);
     auto fields = builder.beginStruct(IdentTy);
     fields.addInt(CGM.Int32Ty, 0);
     fields.addInt(CGM.Int32Ty, Flags);
     fields.addInt(CGM.Int32Ty, 0);
     fields.addInt(CGM.Int32Ty, 0);
void CGOpenMPRuntime::createOffloadEntry(llvm::Constant *ID,
                                          int32_t Flags) {
   StringRef Name = Addr->getName();
   auto *TgtOffloadEntryType = cast<llvm::StructType>(
       CGM.getTypes().ConvertTypeForMem(getTgtOffloadEntryQTy()));
   llvm::LLVMContext &C = CGM.getModule().getContext();
   llvm::Module &M = CGM.getModule();
 
   // Make sure the address has the right type.
-  llvm::Constant *AddrPtr = llvm::ConstantExpr::getBitCast(ID, CGM.VoidPtrTy);
+  llvm::Constant *AddrPtr = llvm::ConstantExpr::getPointerCast(ID, CGM.VoidPtrTy);
 
   // Create constant string with the name.
   llvm::Constant *StrPtrInit = llvm::ConstantDataArray::getString(C, Name);
 
   llvm::GlobalVariable *Str =
       new llvm::GlobalVariable(M, StrPtrInit->getType(), /*isConstant=*/true,
                                llvm::GlobalValue::InternalLinkage, StrPtrInit,
                                ".omp_offloading.entry_name");
   Str->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
-  llvm::Constant *StrPtr = llvm::ConstantExpr::getBitCast(Str, CGM.Int8PtrTy);
+  llvm::Constant *StrPtr = llvm::ConstantExpr::getPointerCast(Str, CGM.Int8PtrTy);
 
   // We can't have any padding between symbols, so we need to have 1-byte
   // alignment.
   auto Align = CharUnits::fromQuantity(1);
 
   // Create the entry struct.
   ConstantInitBuilder EntryBuilder(CGM);
   auto EntryInit = EntryBuilder.beginStruct(TgtOffloadEntryType);
void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
   // need to keep that, and could therefore inline the host function if proven
   // worthwhile during optimization. In the other hand, if emitting code for the
   // device, the ID has to be the function address so that it can retrieved from
   // the offloading entry and launched by the runtime library. We also mark the
   // outlined function to have external linkage in case we are emitting code for
   // the device, because these functions will be entry points to the device.
 
   if (CGM.getLangOpts().OpenMPIsDevice) {
-    OutlinedFnID = llvm::ConstantExpr::getBitCast(OutlinedFn, CGM.Int8PtrTy);
+    OutlinedFnID = llvm::ConstantExpr::getPointerCast(OutlinedFn, CGM.Int8PtrTy);
     OutlinedFn->setLinkage(llvm::GlobalValue::ExternalLinkage);
   } else
     OutlinedFnID = new llvm::GlobalVariable(
         CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
         llvm::GlobalValue::PrivateLinkage,
         llvm::Constant::getNullValue(CGM.Int8Ty), ".omp_offload.region_id");
 
   // Register the information for the entry associated with this target region.

lib/CodeGen/CGVTT.cpp

llvm::Value *Idxs[] = {
        llvm::ConstantInt::get(Int32Ty, AddressPoint.VTableIndex),
        llvm::ConstantInt::get(Int32Ty, AddressPoint.AddressPointIndex),
      };
 
      llvm::Constant *Init = llvm::ConstantExpr::getGetElementPtr(
          VTable->getValueType(), VTable, Idxs, /*InBounds=*/true,
          /*InRangeIndex=*/1);
 
-     Init = llvm::ConstantExpr::getBitCast(Init, Int8PtrTy);
+     Init = llvm::ConstantExpr::getPointerCast(Init, Int8PtrTy);
 
      VTTComponents.push_back(Init);
   }
 
   llvm::Constant *Init = llvm::ConstantArray::get(ArrayType, VTTComponents);
 
   VTT->setInitializer(Init);
 

lib/CodeGen/CGVTables.cpp

void CodeGenVTables::addVTableComponent(
 
   case VTableComponent::CK_VBaseOffset:
     return addOffsetConstant(component.getVBaseOffset());
 
   case VTableComponent::CK_OffsetToTop:
     return addOffsetConstant(component.getOffsetToTop());
 
   case VTableComponent::CK_RTTI:
-    return builder.add(llvm::ConstantExpr::getBitCast(rtti, CGM.Int8PtrTy));
+    return builder.add(llvm::ConstantExpr::getPointerCast(rtti, CGM.Int8PtrTy));
 
   case VTableComponent::CK_FunctionPointer:
   case VTableComponent::CK_CompleteDtorPointer:
   case VTableComponent::CK_DeletingDtorPointer: {
     GlobalDecl GD;
 
     // Get the right global decl.
     switch (component.getKind()) {
case VTableComponent::CK_DeletingDtorPointer: {
     }
 
     auto getSpecialVirtualFn = [&](StringRef name) {
       llvm::FunctionType *fnTy =
           llvm::FunctionType::get(CGM.VoidTy, /*isVarArg=*/false);
       llvm::Constant *fn = CGM.CreateRuntimeFunction(fnTy, name);
       if (auto f = dyn_cast<llvm::Function>(fn))
         f->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
-      return llvm::ConstantExpr::getBitCast(fn, CGM.Int8PtrTy);
+      return llvm::ConstantExpr::getPointerCast(fn, CGM.Int8PtrTy);
     };
 
     llvm::Constant *fnPtr;
 
     // Pure virtual member functions.
     if (cast<CXXMethodDecl>(GD.getDecl())->isPure()) {
       if (!PureVirtualFn)
         PureVirtualFn =
// Thunks.
       fnPtr = CGM.GetAddrOfThunk(GD, thunkInfo);
 
     // Otherwise we can use the method definition directly.
     } else {
       llvm::Type *fnTy = CGM.getTypes().GetFunctionTypeForVTable(GD);
       fnPtr = CGM.GetAddrOfFunction(GD, fnTy, /*ForVTable=*/true);
     }
 
-    fnPtr = llvm::ConstantExpr::getBitCast(fnPtr, CGM.Int8PtrTy);
+    fnPtr = llvm::ConstantExpr::getPointerCast(fnPtr, CGM.Int8PtrTy);
     builder.add(fnPtr);
     return;
   }
 
   case VTableComponent::CK_UnusedFunctionPointer:
     return builder.addNullPointer(CGM.Int8PtrTy);
   }
 

lib/CodeGen/CodeGenFunction.h

struct LifetimeExtendedCleanupHeader {
     /// The kind of cleanup to push: a value from the CleanupKind enumeration.
     CleanupKind Kind;
 
     size_t getSize() const { return Size; }
     CleanupKind getKind() const { return Kind; }
   };
 
   /// i32s containing the indexes of the cleanup destinations.
-  llvm::AllocaInst *NormalCleanupDest;
+  llvm::Instruction *NormalCleanupDest;
 
   unsigned NextCleanupDestIndex;
 
   /// FirstBlockInfo - The head of a singly-linked-list of block layouts.
   CGBlockInfo *FirstBlockInfo;
 
   /// EHResumeBlock - Unified block containing a call to llvm.eh.resume.
   llvm::BasicBlock *EHResumeBlock;
 
   /// The exception slot.  All landing pads write the current exception pointer
   /// into this alloca.
   llvm::Value *ExceptionSlot;
 
   /// The selector slot.  Under the MandatoryCleanup model, all landing pads
-  /// write the current selector value into this alloca.
-  llvm::AllocaInst *EHSelectorSlot;
+  /// write the current selector value into this instruction.
+  llvm::Instruction *EHSelectorSlot;
 
   /// A stack of exception code slots. Entering an __except block pushes a slot
   /// on the stack and leaving pops one. The __exception_code() intrinsic loads
   /// a value from the top of the stack.
   SmallVector<Address, 1> SEHCodeSlotStack;
 
   /// Value returned by __exception_info intrinsic.
   llvm::Value *SEHInfo = nullptr;
class FinallyInfo {
     /// Where the catchall's edge through the cleanup should go.
     JumpDest RethrowDest;
 
     /// A function to call to enter the catch.
     llvm::Constant *BeginCatchFn;
 
     /// An i1 variable indicating whether or not the @finally is
     /// running for an exception.
-    llvm::AllocaInst *ForEHVar;
+    llvm::Instruction *ForEHVar;
 
     /// An i8* variable into which the exception pointer to rethrow
     /// has been saved.
-    llvm::AllocaInst *SavedExnVar;
+    llvm::Instruction *SavedExnVar;
 
   public:
     void enter(CodeGenFunction &CGF, const Stmt *Finally,
                llvm::Constant *beginCatchFn, llvm::Constant *endCatchFn,
                llvm::Constant *rethrowFn);
     void exit(CodeGenFunction &CGF);
   };
 
#endif
   Address EmitLoadOfReference(Address Ref, const ReferenceType *RefTy,
                               AlignmentSource *Source = nullptr);
   LValue EmitLoadOfReferenceLValue(Address Ref, const ReferenceType *RefTy);
 
   Address EmitLoadOfPointer(Address Ptr, const PointerType *PtrTy,
                             AlignmentSource *Source = nullptr);
   LValue EmitLoadOfPointerLValue(Address Ptr, const PointerType *PtrTy);
 
+  /// Create an alloca instruction. If the default address space is not 0,
+  /// insert addrspacecast instruction which casts the alloca instruction
+  /// to the default address space.
+  llvm::Instruction *CreateAlloca(llvm::Type *Ty, const Twine &Name = "tmp",
+                                  llvm::Instruction *InsertPos = nullptr);
   /// CreateTempAlloca - This creates a alloca and inserts it into the entry
   /// block. The caller is responsible for setting an appropriate alignment on
-  /// the alloca.
-  llvm::AllocaInst *CreateTempAlloca(llvm::Type *Ty,
+  /// the alloca. If the default address space is not 0, insert addrspacecast.
+  llvm::Instruction *CreateTempAlloca(llvm::Type *Ty,
                                       const Twine &Name = "tmp");
   Address CreateTempAlloca(llvm::Type *Ty, CharUnits align,
                            const Twine &Name = "tmp");
 
+  /// Get alloca instruction operand of an addrspacecast instruction.
+  /// If \p Inst is alloca instruction, returns \p Inst;
+  llvm::AllocaInst *getAddrSpaceCastedAlloca(llvm::Instruction *Inst) const;
+
   /// CreateDefaultAlignedTempAlloca - This creates an alloca with the
   /// default ABI alignment of the given LLVM type.
   ///
   /// IMPORTANT NOTE: This is *not* generally the right alignment for
   /// any given AST type that happens to have been lowered to the
   /// given IR type.  This should only ever be used for function-local,
   /// IR-driven manipulations like saving and restoring a value.  Do
   /// not hand this address off to arbitrary IRGen routines, and especially

lib/CodeGen/CodeGenFunction.cpp

void CodeGenFunction::EmitFunctionInstrumentation(const char *Fn) {
 
   llvm::Constant *F = CGM.CreateRuntimeFunction(FunctionTy, Fn);
   llvm::CallInst *CallSite = Builder.CreateCall(
     CGM.getIntrinsic(llvm::Intrinsic::returnaddress),
     llvm::ConstantInt::get(Int32Ty, 0),
     "callsite");
 
   llvm::Value *args[] = {
-    llvm::ConstantExpr::getBitCast(CurFn, PointerTy),
+    llvm::ConstantExpr::getPointerCast(CurFn, PointerTy),
     CallSite
   };
 
   EmitNounwindRuntimeCall(F, args);
 }
 
 static void removeImageAccessQualifier(std::string& TyName) {
   std::string ReadOnlyQual("__read_only");

lib/CodeGen/CodeGenModule.cpp

 #include "clang/Basic/SourceManager.h"
 #include "clang/Basic/TargetInfo.h"
 #include "clang/Basic/Version.h"
 #include "clang/Frontend/CodeGenOptions.h"
 #include "clang/Sema/SemaDiagnostic.h"
 #include "llvm/ADT/Triple.h"
 #include "llvm/IR/CallSite.h"
 #include "llvm/IR/CallingConv.h"
+#include "llvm/IR/Constants.h"
 #include "llvm/IR/DataLayout.h"
 #include "llvm/IR/Intrinsics.h"
 #include "llvm/IR/LLVMContext.h"
 #include "llvm/IR/Module.h"
 #include "llvm/ProfileData/InstrProfReader.h"
 #include "llvm/Support/ConvertUTF.h"
 #include "llvm/Support/ErrorHandling.h"
 #include "llvm/Support/MD5.h"
CodeGenModule::CodeGenModule(ASTContext &C, const HeaderSearchOptions &HSO,
                              DiagnosticsEngine &diags,
                              CoverageSourceInfo *CoverageInfo)
     : Context(C), LangOpts(C.getLangOpts()), HeaderSearchOpts(HSO),
       PreprocessorOpts(PPO), CodeGenOpts(CGO), TheModule(M), Diags(diags),
       Target(C.getTargetInfo()), ABI(createCXXABI(*this)),
       VMContext(M.getContext()), Types(*this), VTables(*this),
       SanitizerMD(new SanitizerMetadata(*this)) {
 
+  unsigned DefaultTargetAddressSpace = Target.getDefaultTargetAddressSpace(LangOpts);
   // Initialize the type cache.
   llvm::LLVMContext &LLVMContext = M.getContext();
   VoidTy = llvm::Type::getVoidTy(LLVMContext);
   Int8Ty = llvm::Type::getInt8Ty(LLVMContext);
   Int16Ty = llvm::Type::getInt16Ty(LLVMContext);
   Int32Ty = llvm::Type::getInt32Ty(LLVMContext);
   Int64Ty = llvm::Type::getInt64Ty(LLVMContext);
   FloatTy = llvm::Type::getFloatTy(LLVMContext);
   DoubleTy = llvm::Type::getDoubleTy(LLVMContext);
-  PointerWidthInBits = C.getTargetInfo().getPointerWidth(0);
+  PointerWidthInBits = C.getTargetInfo().getPointerWidth(DefaultTargetAddressSpace);
   PointerAlignInBytes =
-    C.toCharUnitsFromBits(C.getTargetInfo().getPointerAlign(0)).getQuantity();
+    C.toCharUnitsFromBits(C.getTargetInfo().getPointerAlign(DefaultTargetAddressSpace)).getQuantity();
   SizeSizeInBytes =
     C.toCharUnitsFromBits(C.getTargetInfo().getMaxPointerWidth()).getQuantity();
   IntAlignInBytes =
     C.toCharUnitsFromBits(C.getTargetInfo().getIntAlign()).getQuantity();
   IntTy = llvm::IntegerType::get(LLVMContext, C.getTargetInfo().getIntWidth());
   IntPtrTy = llvm::IntegerType::get(LLVMContext,
     C.getTargetInfo().getMaxPointerWidth());
-  Int8PtrTy = Int8Ty->getPointerTo(0);
-  Int8PtrPtrTy = Int8PtrTy->getPointerTo(0);
+  Int8PtrTy = Int8Ty->getPointerTo(DefaultTargetAddressSpace);
+  Int8PtrPtrTy = Int8PtrTy->getPointerTo(DefaultTargetAddressSpace);
 
   RuntimeCC = getTargetCodeGenInfo().getABIInfo().getRuntimeCC();
   BuiltinCC = getTargetCodeGenInfo().getABIInfo().getBuiltinCC();
 
   if (LangOpts.ObjC1)
     createObjCRuntime();
   if (LangOpts.OpenCL)
     createOpenCLRuntime();
void CodeGenModule::EmitCtorList(CtorList &Fns, const char *GlobalName) {
   // Construct the constructor and destructor arrays.
   ConstantInitBuilder builder(*this);
   auto ctors = builder.beginArray(CtorStructTy);
   for (const auto &I : Fns) {
     auto ctor = ctors.beginStruct(CtorStructTy);
     ctor.addInt(Int32Ty, I.Priority);
     ctor.add(llvm::ConstantExpr::getBitCast(I.Initializer, CtorPFTy));
     if (I.AssociatedData)
-      ctor.add(llvm::ConstantExpr::getBitCast(I.AssociatedData, VoidPtrTy));
+      ctor.add(llvm::ConstantExpr::getPointerCast(I.AssociatedData, VoidPtrTy));
     else
       ctor.addNullPointer(VoidPtrTy);
     ctor.finishAndAddTo(ctors);
   }
 
   auto list =
     ctors.finishAndCreateGlobal(GlobalName, getPointerAlign(),
                                 /*constant*/ false,
llvm::Constant *CodeGenModule::EmitAnnotateAttr(llvm::GlobalValue *GV,
                                                 const AnnotateAttr *AA,
                                                 SourceLocation L) {
   // Get the globals for file name, annotation, and the line number.
   llvm::Constant *AnnoGV = EmitAnnotationString(AA->getAnnotation()),
                  *UnitGV = EmitAnnotationUnit(L),
                  *LineNoCst = EmitAnnotationLineNo(L);
 
   // Create the ConstantStruct for the global annotation.
+  unsigned AS = GV->getType()->getAddressSpace();
+  llvm::PointerType *I8PTy = (AS == Int8PtrTy->getAddressSpace()) ?
+    Int8PtrTy : Int8Ty->getPointerTo(AS);
   llvm::Constant *Fields[4] = {
-    llvm::ConstantExpr::getBitCast(GV, Int8PtrTy),
-    llvm::ConstantExpr::getBitCast(AnnoGV, Int8PtrTy),
-    llvm::ConstantExpr::getBitCast(UnitGV, Int8PtrTy),
+    llvm::ConstantExpr::getPointerCast(GV, I8PTy),
+    llvm::ConstantExpr::getPointerCast(AnnoGV, I8PTy),
+    llvm::ConstantExpr::getPointerCast(UnitGV, I8PTy),
     LineNoCst
   };
   return llvm::ConstantStruct::getAnon(Fields);
 }
 
 void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D,
                                          llvm::GlobalValue *GV) {
   assert(D->hasAttr<AnnotateAttr>() && "no annotate attribute");
ConstantAddress CodeGenModule::GetWeakRefReference(const ValueDecl *VD) {
 
   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());
-    auto Ptr = llvm::ConstantExpr::getBitCast(Entry, DeclTy->getPointerTo(AS));
+    auto Ptr = llvm::ConstantExpr::getPointerCast(Entry, DeclTy->getPointerTo(AS));
     return ConstantAddress(Ptr, Alignment);
   }
 
   llvm::Constant *Aliasee;
   if (isa<llvm::FunctionType>(DeclTy))
     Aliasee = GetOrCreateLLVMFunction(AA->getAliasee(), DeclTy,
                                       GlobalDecl(cast<FunctionDecl>(VD)),
                                       /*ForVTable=*/false);

 }
 
 static void ReplaceUsesOfNonProtoTypeWithRealFunction(llvm::GlobalValue *Old,
                                                       llvm::Function *NewFn);
 
 /// GetOrCreateLLVMFunction - If the specified mangled name is not in the
 /// module, create and return an llvm Function with the specified type. If there
 /// is something in the module with the specified name, return it potentially
-/// bitcasted to the right type.
+/// casted to the right type.
 ///
 /// If D is non-null, it specifies a decl that correspond to this.  This is used
 /// to set the attributes on the function when it is first created.
 llvm::Constant *
 CodeGenModule::GetOrCreateLLVMFunction(StringRef MangledName,
                                        llvm::Type *Ty,
                                        GlobalDecl GD, bool ForVTable,
                                        bool DontDefer, bool IsThunk,
if ((isa<llvm::Function>(Entry) || isa<llvm::GlobalAlias>(Entry)) &&
         (Entry->getType()->getElementType() == Ty)) {
       return Entry;
     }
 
     // Make sure the result is of the correct type.
     // (If function is requested for a definition, we always need to create a new
     // function, not just return a bitcast.)
     if (!IsForDefinition)
-      return llvm::ConstantExpr::getBitCast(Entry, Ty->getPointerTo());
+      return llvm::ConstantExpr::getPointerCast(Entry, Ty->getPointerTo());
   }
 
   // This function doesn't have a complete type (for example, the return
   // type is an incomplete struct). Use a fake type instead, and make
   // sure not to try to set attributes.
   bool IsIncompleteFunction = false;
 
   llvm::FunctionType *FTy;
CodeGenModule::GetOrCreateLLVMFunction(StringRef MangledName,
 
   // Make sure the result is of the requested type.
   if (!IsIncompleteFunction) {
     assert(F->getType()->getElementType() == Ty);
     return F;
   }
 
   llvm::Type *PTy = llvm::PointerType::getUnqual(Ty);
-  return llvm::ConstantExpr::getBitCast(F, PTy);
+  return llvm::ConstantExpr::getPointerCast(F, PTy);
 }
 
 /// GetAddrOfFunction - Return the address of the given function.  If Ty is
 /// non-null, then this function will use the specified type if it has to
 /// create it (this occurs when we see a definition of the function).
 llvm::Constant *CodeGenModule::GetAddrOfFunction(GlobalDecl GD,
                                                  llvm::Type *Ty,
                                                  bool ForVTable,
bool CodeGenModule::isTypeConstant(QualType Ty, bool ExcludeCtor) {
   }
 
   return true;
 }
 
 /// GetOrCreateLLVMGlobal - If the specified mangled name is not in the module,
 /// create and return an llvm GlobalVariable with the specified type.  If there
 /// is something in the module with the specified name, return it potentially
-/// bitcasted to the right type.
+/// casted to the right type.
 ///
 /// If D is non-null, it specifies a decl that correspond to this.  This is used
 /// to set the attributes on the global when it is first created.
 ///
 /// If IsForDefinition is true, it is guranteed that an actual global with
 /// type Ty will be returned, not conversion of a variable with the same
 /// mangled name but some other type.
 llvm::Constant *
if (IsForDefinition && !Entry->isDeclaration()) {
           DiagnosedConflictingDefinitions.insert(D).second) {
         getDiags().Report(D->getLocation(),
                           diag::err_duplicate_mangled_name);
         getDiags().Report(OtherGD.getDecl()->getLocation(),
                           diag::note_previous_definition);
       }
     }
 
-    // Make sure the result is of the correct type.
-    if (Entry->getType()->getAddressSpace() != Ty->getAddressSpace())
-      return llvm::ConstantExpr::getAddrSpaceCast(Entry, Ty);
-
     // (If global is requested for a definition, we always need to create a new
     // global, not just return a bitcast.)
     if (!IsForDefinition)
-      return llvm::ConstantExpr::getBitCast(Entry, Ty);
+      return llvm::ConstantExpr::getPointerCast(Entry, Ty);
   }
 
   unsigned AddrSpace = GetGlobalVarAddressSpace(D, Ty->getAddressSpace());
   auto *GV = new llvm::GlobalVariable(
       getModule(), Ty->getElementType(), false,
       llvm::GlobalValue::ExternalLinkage, nullptr, MangledName, nullptr,
       llvm::GlobalVariable::NotThreadLocal, AddrSpace);
 
   // If we already created a global with the same mangled name (but different
   // type) before, take its name and remove it from its parent.
   if (Entry) {
     GV->takeName(Entry);
 
     if (!Entry->use_empty()) {
       llvm::Constant *NewPtrForOldDecl =
-          llvm::ConstantExpr::getBitCast(GV, Entry->getType());
+          llvm::ConstantExpr::getPointerCast(GV, Entry->getType());
       Entry->replaceAllUsesWith(NewPtrForOldDecl);
     }
 
     Entry->eraseFromParent();
   }
 
   // This is the first use or definition of a mangled name.  If there is a
   // deferred decl with this name, remember that we need to emit it at the end
GV = new llvm::GlobalVariable(getModule(), Ty, /*isConstant=*/true,
                                 Linkage, nullptr, Name);
 
   if (OldGV) {
     // Replace occurrences of the old variable if needed.
     GV->takeName(OldGV);
     
     if (!OldGV->use_empty()) {
       llvm::Constant *NewPtrForOldDecl =
-      llvm::ConstantExpr::getBitCast(GV, OldGV->getType());
+      llvm::ConstantExpr::getPointerCast(GV, OldGV->getType());
       OldGV->replaceAllUsesWith(NewPtrForOldDecl);
     }
     
     OldGV->eraseFromParent();
   }
 
   if (supportsCOMDAT() && GV->isWeakForLinker() &&
       !GV->hasAvailableExternallyLinkage())
unsigned CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D,
                                                  unsigned AddrSpace) {
   if (D && LangOpts.CUDA && LangOpts.CUDAIsDevice) {
     if (D->hasAttr<CUDAConstantAttr>())
       AddrSpace = getContext().getTargetAddressSpace(LangAS::cuda_constant);
     else if (D->hasAttr<CUDASharedAttr>())
       AddrSpace = getContext().getTargetAddressSpace(LangAS::cuda_shared);
     else
       AddrSpace = getContext().getTargetAddressSpace(LangAS::cuda_device);
+  } else if (getTriple().getArch() == llvm::Triple::amdgcn &&
+      (LangOpts.CPlusPlus || LangOpts.OpenMP)) {
+    if (D && D->getType().isConstant(getContext()))
+      AddrSpace = getContext().getTargetAddressSpace(LangAS::opencl_constant);
+    else
+      AddrSpace = getContext().getTargetAddressSpace(LangAS::opencl_global);
   }
 
   return AddrSpace;
 }
 
 template<typename SomeDecl>
 void CodeGenModule::MaybeHandleStaticInExternC(const SomeDecl *D,
                                                llvm::GlobalValue *GV) {
if (!GV ||
     Entry->setName(StringRef());
 
     // Make a new global with the correct type, this is now guaranteed to work.
     GV = cast<llvm::GlobalVariable>(
         GetAddrOfGlobalVar(D, InitType, ForDefinition_t(!IsTentative)));
 
     // Replace all uses of the old global with the new global
     llvm::Constant *NewPtrForOldDecl =
-        llvm::ConstantExpr::getBitCast(GV, Entry->getType());
+        llvm::ConstantExpr::getPointerCast(GV, Entry->getType());
     Entry->replaceAllUsesWith(NewPtrForOldDecl);
 
     // Erase the old global, since it is no longer used.
     cast<llvm::GlobalValue>(Entry)->eraseFromParent();
   }
 
   MaybeHandleStaticInExternC(D, GV);
 
if (Entry) {
     // by the alias, as in:
     //   extern int test6();
     //   ...
     //   int test6() __attribute__((alias("test7")));
     //
     // Remove it and replace uses of it with the alias.
     GA->takeName(Entry);
 
-    Entry->replaceAllUsesWith(llvm::ConstantExpr::getBitCast(GA,
+    Entry->replaceAllUsesWith(llvm::ConstantExpr::getPointerCast(GA,
                                                           Entry->getType()));
     Entry->eraseFromParent();
   } else {
     GA->setName(MangledName);
   }
 
   // Set attributes which are particular to an alias; this is a
   // specialization of the attributes which may be set on a global
GV->setSection(isUTF16 ? "__TEXT,__ustring"
                            : "__TEXT,__cstring,cstring_literals");
 
   // String.
   llvm::Constant *Str =
       llvm::ConstantExpr::getGetElementPtr(GV->getValueType(), GV, Zeros);
 
   if (isUTF16)
     // Cast the UTF16 string to the correct type.
-    Str = llvm::ConstantExpr::getBitCast(Str, Int8PtrTy);
+    Str = llvm::ConstantExpr::getPointerCast(Str, Int8PtrTy);
   Fields.add(Str);
 
   // String length.
   auto Ty = getTypes().ConvertType(getContext().LongTy);
   Fields.addInt(cast<llvm::IntegerType>(Ty), StringLength);
 
   CharUnits Alignment = getPointerAlign();
 
CodeGenModule::GetConstantArrayFromStringLiteral(const StringLiteral *E) {
   return llvm::ConstantDataArray::get(VMContext, Elements);
 }
 
 static llvm::GlobalVariable *
 GenerateStringLiteral(llvm::Constant *C, llvm::GlobalValue::LinkageTypes LT,
                       CodeGenModule &CGM, StringRef GlobalName,
                       CharUnits Alignment) {
   // OpenCL v1.2 s6.5.3: a string literal is in the constant address space.
-  unsigned AddrSpace = 0;
+  unsigned AddrSpace = CGM.getContext().getTargetConstantAddressSpace();
   if (CGM.getLangOpts().OpenCL)
     AddrSpace = CGM.getContext().getTargetAddressSpace(LangAS::opencl_constant);
 
   llvm::Module &M = CGM.getModule();
   // Create a global variable for this string
   auto *GV = new llvm::GlobalVariable(
       M, C->getType(), !CGM.getLangOpts().WritableStrings, LT, C, GlobalName,
       nullptr, llvm::GlobalVariable::NotThreadLocal, AddrSpace);
for (auto *I : DC->decls()) {
     }
 
     EmitTopLevelDecl(I);
   }
 }
 
 /// EmitTopLevelDecl - Emit code for a single top level declaration.
 void CodeGenModule::EmitTopLevelDecl(Decl *D) {
+  if (getenv("DBG_CG_DECL")) {
+    llvm::errs() << "decl: "; D->dump();
+  }
   // Ignore dependent declarations.
   if (D->getDeclContext() && D->getDeclContext()->isDependentContext())
     return;
 
   switch (D->getKind()) {
   case Decl::CXXConversion:
   case Decl::CXXMethod:
   case Decl::Function:

lib/CodeGen/CodeGenTypes.h

public:
   const ABIInfo &getABIInfo() const { return TheABIInfo; }
   const TargetInfo &getTarget() const { return Target; }
   CGCXXABI &getCXXABI() const { return TheCXXABI; }
   llvm::LLVMContext &getLLVMContext() { return TheModule.getContext(); }
 
   /// ConvertType - Convert type T into a llvm::Type.
   llvm::Type *ConvertType(QualType T);
 
+  /// Get a pointer type pointing to the given QualType \p T.
+  llvm::PointerType *getPointerTypeTo(QualType T = QualType());
+
+  /// Get a pointer type pointing to the given llvm::Type \p T in the default
+  /// target address space.
+  llvm::PointerType *getDefaultPointerTo(llvm::Type *T);
+
   /// \brief Converts the GlobalDecl into an llvm::Type. This should be used
   /// when we know the target of the function we want to convert.  This is
   /// because some functions (explicitly, those with pass_object_size
   /// parameters) may not have the same signature as their type portrays, and
   /// can only be called directly.
   llvm::Type *ConvertFunctionType(QualType FT,
                                   const FunctionDecl *FD = nullptr);
 

lib/CodeGen/CodeGenTypes.cpp

if (SkippedLayout)
     TypeCache.clear();
 
   if (RecordsBeingLaidOut.empty())
     while (!DeferredRecords.empty())
       ConvertRecordDeclType(DeferredRecords.pop_back_val());
   return ResultType;
 }
 
+llvm::PointerType *CodeGenTypes::getPointerTypeTo(QualType T) {
+  return ConvertType(T)->getPointerTo(Context.getTargetAddressSpace(T));
+}
+
+llvm::PointerType *CodeGenTypes::getDefaultPointerTo(llvm::Type *T) {
+  return T->getPointerTo(Context.getTargetDefaultAddressSpace());
+}
+
 /// ConvertType - Convert the specified type to its LLVM form.
 llvm::Type *CodeGenTypes::ConvertType(QualType T) {
   T = Context.getCanonicalType(T);
 
   const Type *Ty = T.getTypePtr();
 
   // RecordTypes are cached and processed specially.
   if (const RecordType *RT = dyn_cast<RecordType>(Ty))

lib/CodeGen/ItaniumCXXABI.cpp

   // The address of the destructor.  If the exception type has a
   // trivial destructor (or isn't a record), we just pass null.
   llvm::Constant *Dtor = nullptr;
   if (const RecordType *RecordTy = ThrowType->getAs<RecordType>()) {
     CXXRecordDecl *Record = cast<CXXRecordDecl>(RecordTy->getDecl());
     if (!Record->hasTrivialDestructor()) {
       CXXDestructorDecl *DtorD = Record->getDestructor();
       Dtor = CGM.getAddrOfCXXStructor(DtorD, StructorType::Complete);
-      Dtor = llvm::ConstantExpr::getBitCast(Dtor, CGM.Int8PtrTy);
+      Dtor = llvm::ConstantExpr::getPointerCast(Dtor, CGM.Int8PtrTy);
     }
   }
   if (!Dtor) Dtor = llvm::Constant::getNullValue(CGM.Int8PtrTy);
 
   llvm::Value *args[] = { ExceptionPtr, TypeInfo, Dtor };
   CGF.EmitNoreturnRuntimeCallOrInvoke(getThrowFn(CGM), args);
 }
 

 
 llvm::Value *ItaniumCXXABI::EmitTypeid(CodeGenFunction &CGF,
                                        QualType SrcRecordTy,
                                        Address ThisPtr,
                                        llvm::Type *StdTypeInfoPtrTy) {
   auto *ClassDecl =
       cast<CXXRecordDecl>(SrcRecordTy->getAs<RecordType>()->getDecl());
   llvm::Value *Value =
-      CGF.GetVTablePtr(ThisPtr, StdTypeInfoPtrTy->getPointerTo(), ClassDecl);
+      CGF.GetVTablePtr(ThisPtr, CGF.getTypes().getDefaultPointerTo(
+          StdTypeInfoPtrTy), ClassDecl);
 
   // Load the type info.
   Value = CGF.Builder.CreateConstInBoundsGEP1_64(Value, -1ULL);
   return CGF.Builder.CreateAlignedLoad(Value, CGF.getPointerAlign());
 }
 
 bool ItaniumCXXABI::shouldDynamicCastCallBeNullChecked(bool SrcIsPtr,
                                                        QualType SrcRecordTy) {

       guardTy = CGF.SizeTy;
       guardAlignment = CGF.getSizeAlign();
     } else {
       guardTy = CGF.Int64Ty;
       guardAlignment = CharUnits::fromQuantity(
                              CGM.getDataLayout().getABITypeAlignment(guardTy));
     }
   }
-  llvm::PointerType *guardPtrTy = guardTy->getPointerTo();
+  llvm::PointerType *guardPtrTy = guardTy->getPointerTo(
+      getContext().getTargetDefaultAddressSpace());
 
   // Create the guard variable if we don't already have it (as we
   // might if we're double-emitting this function body).
   llvm::GlobalVariable *guard = CGM.getStaticLocalDeclGuardAddress(&D);
   if (!guard) {
     // Mangle the name for the guard.
     SmallString<256> guardName;
     {
       llvm::raw_svector_ostream out(guardName);
       getMangleContext().mangleStaticGuardVariable(&D, out);
     }
 
     // Create the guard variable with a zero-initializer.
     // Just absorb linkage and visibility from the guarded variable.
     guard = new llvm::GlobalVariable(CGM.getModule(), guardTy,
                                      false, var->getLinkage(),
                                      llvm::ConstantInt::get(guardTy, 0),
-                                     guardName.str());
+                                     guardName.str(),
+                                     /* InsertBefore */ nullptr,
+                                     llvm::GlobalValue::NotThreadLocal,
+                                     getContext().getTargetGlobalAddressSpace());
     guard->setVisibility(var->getVisibility());
     // If the variable is thread-local, so is its guard variable.
     guard->setThreadLocalMode(var->getThreadLocalMode());
     guard->setAlignment(guardAlignment.getQuantity());
 
     // The ABI says: "It is suggested that it be emitted in the same COMDAT
     // group as the associated data object." In practice, this doesn't work for
     // non-ELF and non-Wasm object formats, so only do it for ELF and Wasm.

   // Create a variable that binds the atexit to this shared object.
   llvm::Constant *handle =
       CGF.CGM.CreateRuntimeVariable(CGF.Int8Ty, "__dso_handle");
   auto *GV = cast<llvm::GlobalValue>(handle->stripPointerCasts());
   GV->setVisibility(llvm::GlobalValue::HiddenVisibility);
 
   llvm::Value *args[] = {
     llvm::ConstantExpr::getBitCast(dtor, dtorTy),
-    llvm::ConstantExpr::getBitCast(addr, CGF.Int8PtrTy),
-    handle
+    llvm::ConstantExpr::getPointerCast(addr, CGF.Int8PtrTy),
+    llvm::ConstantExpr::getPointerCast(handle, CGF.Int8PtrTy)
   };
   CGF.EmitNounwindRuntimeCall(atexit, args);
 }
 
 /// Register a global destructor as best as we know how.
 void ItaniumCXXABI::registerGlobalDtor(CodeGenFunction &CGF,
                                        const VarDecl &D,
                                        llvm::Constant *dtor,

                                   Name);
     if (const RecordType *RecordTy = dyn_cast<RecordType>(Ty)) {
       const CXXRecordDecl *RD = cast<CXXRecordDecl>(RecordTy->getDecl());
       if (RD->hasAttr<DLLImportAttr>())
         GV->setDLLStorageClass(llvm::GlobalVariable::DLLImportStorageClass);
     }
   }
 
-  return llvm::ConstantExpr::getBitCast(GV, CGM.Int8PtrTy);
+  return llvm::ConstantExpr::getPointerCast(GV, CGM.Int8PtrTy);
 }
 
 /// 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

 
   llvm::Type *PtrDiffTy =
     CGM.getTypes().ConvertType(CGM.getContext().getPointerDiffType());
 
   // The vtable address point is 2.
   llvm::Constant *Two = llvm::ConstantInt::get(PtrDiffTy, 2);
   VTable =
       llvm::ConstantExpr::getInBoundsGetElementPtr(CGM.Int8PtrTy, VTable, Two);
-  VTable = llvm::ConstantExpr::getBitCast(VTable, CGM.Int8PtrTy);
+  VTable = llvm::ConstantExpr::getPointerCast(VTable, CGM.Int8PtrTy);
 
   Fields.push_back(VTable);
 }
 
 /// \brief 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::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 llvm::ConstantExpr::getPointerCast(OldGV, CGM.Int8PtrTy);
   }
 
   // Check if there is already an external RTTI descriptor for this type.
   bool IsStdLib = IsStandardLibraryRTTIDescriptor(Ty);
   if (!Force && (IsStdLib || ShouldUseExternalRTTIDescriptor(CGM, Ty)))
     return GetAddrOfExternalRTTIDescriptor(Ty);
 
   // Emit the standard library with external linkage.

     // 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);
   } else {
-    TypeNameField = llvm::ConstantExpr::getBitCast(TypeName, CGM.Int8PtrTy);
+    TypeNameField = llvm::ConstantExpr::getPointerCast(TypeName, CGM.Int8PtrTy);
   }
   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:

       // Because the typename and the typeinfo are DLL import, convert them to
       // declarations rather than definitions.  The initializers still need to
       // be constructed to calculate the type for the declarations.
       TypeName->setInitializer(nullptr);
       GV->setInitializer(nullptr);
     }
   }
 
-  return llvm::ConstantExpr::getBitCast(GV, CGM.Int8PtrTy);
+  return llvm::ConstantExpr::getPointerCast(GV, CGM.Int8PtrTy);
 }
 
 /// 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));

test/CodeGenCUDA/address-spaces.cu

-// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck --check-prefixes=NVPTX,CHECK %s
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple amdgcn | FileCheck --check-prefixes=AMDGCN,CHECK %s
 
 // Verifies Clang emits correct address spaces and addrspacecast instructions
 // for CUDA code.
 
 #include "Inputs/cuda.h"
 
 // CHECK: @i = addrspace(1) externally_initialized global
 __device__ int i;
 
-// CHECK: @j = addrspace(4) externally_initialized global
+// AMDGCN: @j = addrspace(2) externally_initialized global
+// NVPTX: @j = addrspace(4) externally_initialized global
 __constant__ int j;
 
 // CHECK: @k = addrspace(3) global
 __shared__ int k;
 
 struct MyStruct {
   int data1;
   int data2;
 };
 
 // CHECK: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct zeroinitializer
 // CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] zeroinitializer
 // CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00
 // CHECK: @b = addrspace(3) global float undef
 
 __device__ void foo() {
-  // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*)
+  // NVPTX: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*)
+  // AMDGCN: load i32, i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @i to i32 addrspace(4)*)
   i++;
 
-  // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @j to i32*)
+  // NVPTX: load i32, i32* addrspacecast (i32 addrspace(4)* @j to i32*)
+  // AMDGCN: load i32, i32 addrspace(4)* addrspacecast (i32 addrspace(2)* @j to i32 addrspace(4)*)
   j++;
 
-  // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*)
+  // NVPTX: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*)
+  // AMDGCN: load i32, i32 addrspace(4)* addrspacecast (i32 addrspace(3)* @k to i32 addrspace(4)*)
   k++;
 
   __shared__ int lk;
-  // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*)
+  // NVPTX: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*)
+  // AMDGCN: load i32, i32 addrspace(4)* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32 addrspace(4)*)
   lk++;
 }
 
 __device__ void func0() {
   __shared__ MyStruct a;
   MyStruct *ap = &a; // composite type
   ap->data1 = 1;
   ap->data2 = 2;
 }
-// CHECK: define void @_Z5func0v()
-// CHECK: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %ap
+// CHECK-LABEL: define void @_Z5func0v()
+// NVPTX: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %ap
+// AMDGCN: store %struct.MyStruct addrspace(4)* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct addrspace(4)*), %struct.MyStruct addrspace(4)* addrspace(4)* %ap
 
 __device__ void callee(float *ap) {
   *ap = 1.0f;
 }
 
 __device__ void func1() {
   __shared__ float a;
   callee(&a); // implicit cast from parameters
 }
-// CHECK: define void @_Z5func1v()
-// CHECK: call void @_Z6calleePf(float* addrspacecast (float addrspace(3)* @_ZZ5func1vE1a to float*))
+// CHECK-LABEL: define void @_Z5func1v()
+// NVPTX: call void @_Z6calleePf(float* addrspacecast (float addrspace(3)* @_ZZ5func1vE1a to float*))
+// AMDGCN: call void @_Z6calleePf(float addrspace(4)* addrspacecast (float addrspace(3)* @_ZZ5func1vE1a to float addrspace(4)*))
 
 __device__ void func2() {
   __shared__ float a[256];
   float *ap = &a[128]; // implicit cast from a decayed array
   *ap = 1.0f;
 }
-// CHECK: define void @_Z5func2v()
-// CHECK: store float* getelementptr inbounds ([256 x float], [256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i32 0, i32 128), float** %ap
-
+// CHECK-LABEL: define void @_Z5func2v()
+// NVPTX: store float* getelementptr inbounds ([256 x float], [256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i32 0, i32 128), float** %ap
+// AMDGCN: store float addrspace(4)* getelementptr inbounds ([256 x float], [256 x float] addrspace(4)* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float] addrspace(4)*), i64 0, i64 128), float addrspace(4)* addrspace(4)* %ap
 __device__ void func3() {
   __shared__ float a;
   float *ap = reinterpret_cast<float *>(&a); // explicit cast
   *ap = 1.0f;
 }
-// CHECK: define void @_Z5func3v()
-// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %ap
+// CHECK-LABEL: define void @_Z5func3v()
+// NVPTX: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %ap
+// AMDGCN: store float addrspace(4)* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float addrspace(4)*), float addrspace(4)* addrspace(4)* %ap
 
 __device__ void func4() {
   __shared__ float a;
   float *ap = (float *)&a; // explicit c-style cast
   *ap = 1.0f;
 }
-// CHECK: define void @_Z5func4v()
-// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %ap
+// CHECK-LABEL: define void @_Z5func4v()
+// NVPTX: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %ap
+// AMDGCN: store float addrspace(4)* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float addrspace(4)*), float addrspace(4)* addrspace(4)* %ap
 
 __shared__ float b;
 
 __device__ float *func5() {
   return &b; // implicit cast from a return value
 }
-// CHECK: define float* @_Z5func5v()
-// CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*)
+// NVPTX-LABEL: define float* @_Z5func5v()
+// AMDGCN-LABEL: define float addrspace(4)* @_Z5func5v()
+// NVPTX: ret float* addrspacecast (float addrspace(3)* @b to float*)
+// AMDGCN: ret float addrspace(4)* addrspacecast (float addrspace(3)* @b to float addrspace(4)*)

test/CodeGenCUDA/convergent.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \
+// RUN:   -disable-llvm-passes -o - %s -DNVPTX | FileCheck -check-prefixes=DEVICE,NVPTX %s
+
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn -emit-llvm \
 // RUN:   -disable-llvm-passes -o - %s | FileCheck -check-prefix DEVICE %s
 
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm \
 // RUN:   -disable-llvm-passes -o - %s | \
 // RUN:  FileCheck -check-prefix HOST %s
 
 #include "Inputs/cuda.h"
 
 // DEVICE: Function Attrs:
 // DEVICE-SAME: convergent
 // DEVICE-NEXT: define void @_Z3foov
 __device__ void foo() {}
 
 // HOST: Function Attrs:
 // HOST-NOT: convergent
 // HOST-NEXT: define void @_Z3barv
 // DEVICE: Function Attrs:
 // DEVICE-SAME: convergent
 // DEVICE-NEXT: define void @_Z3barv
 __host__ __device__ void baz();
 __host__ __device__ void bar() {
   // DEVICE: call void @_Z3bazv() [[CALL_ATTR:#[0-9]+]]
   baz();
-  // DEVICE: call i32 asm "trap;", "=l"() [[ASM_ATTR:#[0-9]+]]
+  #ifdef NVPTX
+  // NVPTX: call i32 asm "trap;", "=l"() [[ASM_ATTR:#[0-9]+]]
   int x;
   asm ("trap;" : "=l"(x));
+  #endif
   // DEVICE: call void asm sideeffect "trap;", ""() [[ASM_ATTR:#[0-9]+]]
   asm volatile ("trap;");
 }
 
 // DEVICE: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
 // DEVICE: attributes [[BAZ_ATTR]] = {
 // DEVICE-SAME: convergent
 // DEVICE-SAME: }
 // DEVICE-DAG: attributes [[CALL_ATTR]] = { convergent
 // DEVICE-DAG: attributes [[ASM_ATTR]] = { convergent
 
 // HOST: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
 // HOST: attributes [[BAZ_ATTR]] = {
 // HOST-NOT: convergent
 // NOST-SAME: }

test/CodeGenCUDA/device-var-init.cu

 // REQUIRES: nvptx-registered-target
 
 // Make sure we don't allow dynamic initialization for device
 // variables, but accept empty constructors allowed by CUDA.
 
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \
-// RUN:     -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck %s
+// RUN:     -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK,NVPTX %s
+
+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 \
+// RUN:     -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck --check-prefixes=CHECK,AMDGCN %s
 
 #ifdef __clang__
 #include "Inputs/cuda.h"
 #endif
 
 // Use the types we share with Sema tests.
 #include "Inputs/cuda-initializers.h"
 
 __device__ int d_v;
 // CHECK: @d_v = addrspace(1) externally_initialized global i32 0,
 __shared__ int s_v;
 // CHECK: @s_v = addrspace(3) global i32 undef,
 __constant__ int c_v;
-// CHECK: addrspace(4) externally_initialized global i32 0,
+// NVPTX: addrspace(4) externally_initialized global i32 0,
+// AMDGCN: addrspace(2) externally_initialized global i32 0,
 
 __device__ int d_v_i = 1;
 // CHECK: @d_v_i = addrspace(1) externally_initialized global i32 1,
 
 // trivial constructor -- allowed
 __device__ T d_t;
 // CHECK: @d_t = addrspace(1) externally_initialized global %struct.T zeroinitializer
 __shared__ T s_t;
 // CHECK: @s_t = addrspace(3) global %struct.T undef,
 __constant__ T c_t;
-// CHECK: @c_t = addrspace(4) externally_initialized global %struct.T zeroinitializer,
+// NVPTX: @c_t = addrspace(4) externally_initialized global %struct.T zeroinitializer,
+// AMDGCN: @c_t = addrspace(2) externally_initialized global %struct.T zeroinitializer,
 
 __device__ T d_t_i = {2};
 // CHECK: @d_t_i = addrspace(1) externally_initialized global %struct.T { i32 2 },
 __constant__ T c_t_i = {2};
-// CHECK: @c_t_i = addrspace(4) externally_initialized global %struct.T { i32 2 },
+// NVPTX: @c_t_i = addrspace(4) externally_initialized global %struct.T { i32 2 },
+// AMDGCN: @c_t_i = addrspace(2) externally_initialized global %struct.T { i32 2 },
 
 // empty constructor
 __device__ EC d_ec;
 // CHECK: @d_ec = addrspace(1) externally_initialized global %struct.EC zeroinitializer,
 __shared__ EC s_ec;
 // CHECK: @s_ec = addrspace(3) global %struct.EC undef,
 __constant__ EC c_ec;
-// CHECK: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer,
+// NVPTX: @c_ec = addrspace(4) externally_initialized global %struct.EC zeroinitializer,
+// AMDGCN: @c_ec = addrspace(2) externally_initialized global %struct.EC zeroinitializer,
 
 // empty destructor
 __device__ ED d_ed;
 // CHECK: @d_ed = addrspace(1) externally_initialized global %struct.ED zeroinitializer,
 __shared__ ED s_ed;
 // CHECK: @s_ed = addrspace(3) global %struct.ED undef,
 __constant__ ED c_ed;
-// CHECK: @c_ed = addrspace(4) externally_initialized global %struct.ED zeroinitializer,
+// NVPTX: @c_ed = addrspace(4) externally_initialized global %struct.ED zeroinitializer,
+// AMDGCN: @c_ed = addrspace(2) externally_initialized global %struct.ED zeroinitializer,
 
 __device__ ECD d_ecd;
 // CHECK: @d_ecd = addrspace(1) externally_initialized global %struct.ECD zeroinitializer,
 __shared__ ECD s_ecd;
 // CHECK: @s_ecd = addrspace(3) global %struct.ECD undef,
 __constant__ ECD c_ecd;
-// CHECK: @c_ecd = addrspace(4) externally_initialized global %struct.ECD zeroinitializer,
+// NVPTX: @c_ecd = addrspace(4) externally_initialized global %struct.ECD zeroinitializer,
+// AMDGCN: @c_ecd = addrspace(2) externally_initialized global %struct.ECD zeroinitializer,
 
 // empty templated constructor -- allowed with no arguments
 __device__ ETC d_etc;
 // CHECK: @d_etc = addrspace(1) externally_initialized global %struct.ETC zeroinitializer,
 __shared__ ETC s_etc;
 // CHECK: @s_etc = addrspace(3) global %struct.ETC undef,
 __constant__ ETC c_etc;
-// CHECK: @c_etc = addrspace(4) externally_initialized global %struct.ETC zeroinitializer,
+// NVPTX: @c_etc = addrspace(4) externally_initialized global %struct.ETC zeroinitializer,
+// AMDGCN: @c_etc = addrspace(2) externally_initialized global %struct.ETC zeroinitializer,
 
 __device__ NCFS d_ncfs;
 // CHECK: @d_ncfs = addrspace(1) externally_initialized global %struct.NCFS { i32 3 }
 __constant__ NCFS c_ncfs;
-// CHECK: @c_ncfs = addrspace(4) externally_initialized global %struct.NCFS { i32 3 }
+// NVPTX: @c_ncfs = addrspace(4) externally_initialized global %struct.NCFS { i32 3 }
+// AMDGCN: @c_ncfs = addrspace(2) externally_initialized global %struct.NCFS { i32 3 }
 
 // Regular base class -- allowed
 __device__ T_B_T d_t_b_t;
 // CHECK: @d_t_b_t = addrspace(1) externally_initialized global %struct.T_B_T zeroinitializer,
 __shared__ T_B_T s_t_b_t;
 // CHECK: @s_t_b_t = addrspace(3) global %struct.T_B_T undef,
 __constant__ T_B_T c_t_b_t;
-// CHECK: @c_t_b_t = addrspace(4) externally_initialized global %struct.T_B_T zeroinitializer,
+// NVPTX: @c_t_b_t = addrspace(4) externally_initialized global %struct.T_B_T zeroinitializer,
+// AMDGCN: @c_t_b_t = addrspace(2) externally_initialized global %struct.T_B_T zeroinitializer,
 
 // Incapsulated object of allowed class -- allowed
 __device__ T_F_T d_t_f_t;
 // CHECK: @d_t_f_t = addrspace(1) externally_initialized global %struct.T_F_T zeroinitializer,
 __shared__ T_F_T s_t_f_t;
 // CHECK: @s_t_f_t = addrspace(3) global %struct.T_F_T undef,
 __constant__ T_F_T c_t_f_t;
-// CHECK: @c_t_f_t = addrspace(4) externally_initialized global %struct.T_F_T zeroinitializer,
+// NVPTX: @c_t_f_t = addrspace(4) externally_initialized global %struct.T_F_T zeroinitializer,
+// AMDGCN: @c_t_f_t = addrspace(2) externally_initialized global %struct.T_F_T zeroinitializer,
 
 // array of allowed objects -- allowed
 __device__ T_FA_T d_t_fa_t;
 // CHECK: @d_t_fa_t = addrspace(1) externally_initialized global %struct.T_FA_T zeroinitializer,
 __shared__ T_FA_T s_t_fa_t;
 // CHECK: @s_t_fa_t = addrspace(3) global %struct.T_FA_T undef,
 __constant__ T_FA_T c_t_fa_t;
-// CHECK: @c_t_fa_t = addrspace(4) externally_initialized global %struct.T_FA_T zeroinitializer,
+// NVPTX: @c_t_fa_t = addrspace(4) externally_initialized global %struct.T_FA_T zeroinitializer,
+// AMDGCN: @c_t_fa_t = addrspace(2) externally_initialized global %struct.T_FA_T zeroinitializer,
 
 
 // Calling empty base class initializer is OK
 __device__ EC_I_EC d_ec_i_ec;
 // CHECK: @d_ec_i_ec = addrspace(1) externally_initialized global %struct.EC_I_EC zeroinitializer,
 __shared__ EC_I_EC s_ec_i_ec;
 // CHECK: @s_ec_i_ec = addrspace(3) global %struct.EC_I_EC undef,
 __constant__ EC_I_EC c_ec_i_ec;
-// CHECK: @c_ec_i_ec = addrspace(4) externally_initialized global %struct.EC_I_EC zeroinitializer,
+// NVPTX: @c_ec_i_ec = addrspace(4) externally_initialized global %struct.EC_I_EC zeroinitializer,
+// AMDGCN: @c_ec_i_ec = addrspace(2) externally_initialized global %struct.EC_I_EC zeroinitializer,
 
 // We should not emit global initializers for device-side variables.
 // CHECK-NOT: @__cxx_global_var_init
 
 // Make sure that initialization restrictions do not apply to local
 // variables.
 __device__ void df() {
   T t;
   // CHECK-NOT: call
   EC ec;
-  // CHECK:   call void @_ZN2ECC1Ev(%struct.EC* %ec)
+  // NVPTX:   call void @_ZN2ECC1Ev(%struct.EC* %ec)
+  // AMDGCN:   call void @_ZN2ECC1Ev(%struct.EC addrspace(4)* %ec)
   ED ed;
   // CHECK-NOT: call
   ECD ecd;
-  // CHECK:   call void @_ZN3ECDC1Ev(%struct.ECD* %ecd)
+  // NVPTX:   call void @_ZN3ECDC1Ev(%struct.ECD* %ecd)
+  // AMDGCN:   call void @_ZN3ECDC1Ev(%struct.ECD addrspace(4)* %ecd)
   ETC etc;
-  // CHECK:   call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc)
+  // NVPTX:   call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc)
+  // AMDGCN:   call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC addrspace(4)* %etc)
   UC uc;
   // undefined constructor -- not allowed
-  // CHECK:   call void @_ZN2UCC1Ev(%struct.UC* %uc)
+  // NVPTX:   call void @_ZN2UCC1Ev(%struct.UC* %uc)
+  // AMDGCN:   call void @_ZN2UCC1Ev(%struct.UC addrspace(4)* %uc)
   UD ud;
   // undefined destructor -- not allowed
   // CHECK-NOT: call
   ECI eci;
   // empty constructor w/ initializer list -- not allowed
-  // CHECK:   call void @_ZN3ECIC1Ev(%struct.ECI* %eci)
+  // NVPTX:   call void @_ZN3ECIC1Ev(%struct.ECI* %eci)
+  // AMDGCN:   call void @_ZN3ECIC1Ev(%struct.ECI addrspace(4)* %eci)
   NEC nec;
   // non-empty constructor -- not allowed
-  // CHECK:   call void @_ZN3NECC1Ev(%struct.NEC* %nec)
+  // NVPTX:   call void @_ZN3NECC1Ev(%struct.NEC* %nec)
+  // AMDGCN:   call void @_ZN3NECC1Ev(%struct.NEC addrspace(4)* %nec)
   // non-empty destructor -- not allowed
   NED ned;
   // no-constructor,  virtual method -- not allowed
-  // CHECK:   call void @_ZN3NCVC1Ev(%struct.NCV* %ncv)
+  // NVPTX:   call void @_ZN3NCVC1Ev(%struct.NCV* %ncv)
+  // AMDGCN:   call void @_ZN3NCVC1Ev(%struct.NCV addrspace(4)* %ncv)
   NCV ncv;
   // CHECK-NOT: call
   VD vd;
-  // CHECK:   call void @_ZN2VDC1Ev(%struct.VD* %vd)
+  // NVPTX:   call void @_ZN2VDC1Ev(%struct.VD* %vd)
+  // AMDGCN:   call void @_ZN2VDC1Ev(%struct.VD addrspace(4)* %vd)
   NCF ncf;
-  // CHECK:   call void @_ZN3NCFC1Ev(%struct.NCF* %ncf)
+  // NVPTX:   call void @_ZN3NCFC1Ev(%struct.NCF* %ncf)
+  // AMDGCN:   call void @_ZN3NCFC1Ev(%struct.NCF addrspace(4)* %ncf)
   NCFS ncfs;
-  // CHECK:   call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs)
+  // NVPTX:   call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs)
+  // AMDGCN:   call void @_ZN4NCFSC1Ev(%struct.NCFS addrspace(4)* %ncfs)
   UTC utc;
-  // CHECK:   call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc)
+  // NVPTX:   call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc)
+  // AMDGCN:   call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC addrspace(4)* %utc)
   NETC netc;
-  // CHECK:   call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc)
+  // NVPTX:   call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc)
+  // AMDGCN:   call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC addrspace(4)* %netc)
   T_B_T t_b_t;
   // CHECK-NOT: call
   T_F_T t_f_t;
   // CHECK-NOT: call
   T_FA_T t_fa_t;
   // CHECK-NOT: call
   EC_I_EC ec_i_ec;
-  // CHECK:   call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec)
+  // NVPTX:   call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec)
+  // AMDGCN:   call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC addrspace(4)* %ec_i_ec)
   EC_I_EC1 ec_i_ec1;
-  // CHECK:   call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1)
+  // NVPTX:   call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1)
+  // AMDGCN:   call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1 addrspace(4)* %ec_i_ec1)
   T_V_T t_v_t;
-  // CHECK:   call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t)
+  // NVPTX:   call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t)
+  // AMDGCN:   call void @_ZN5T_V_TC1Ev(%struct.T_V_T addrspace(4)* %t_v_t)
   T_B_NEC t_b_nec;
-  // CHECK:   call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec)
+  // NVPTX:   call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec)
+  // AMDGCN:   call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC addrspace(4)* %t_b_nec)
   T_F_NEC t_f_nec;
-  // CHECK:   call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
+  // NVPTX:   call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
+  // AMDGCN:   call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC addrspace(4)* %t_f_nec)
   T_FA_NEC t_fa_nec;
-  // CHECK:   call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
+  // NVPTX:   call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
+  // AMDGCN:   call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC addrspace(4)* %t_fa_nec)
   T_B_NED t_b_ned;
   // CHECK-NOT: call
   T_F_NED t_f_ned;
   // CHECK-NOT: call
   T_FA_NED t_fa_ned;
   // CHECK-NOT: call
   static __shared__ EC s_ec;
-  // CHECK-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
+  // NVPTX-NOT: call void @_ZN2ECC1Ev(%struct.EC* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC*))
+  // AMDGCN-NOT: call void @_ZN2ECC1Ev(%struct.EC addrspace(4)* addrspacecast (%struct.EC addrspace(3)* @_ZZ2dfvE4s_ec to %struct.EC addrspace(4)*))
   static __shared__ ETC s_etc;
-  // CHECK-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
+  // NVPTX-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC*))
+  // AMDGCN-NOT: call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC addrspace(4)* addrspacecast (%struct.ETC addrspace(3)* @_ZZ2dfvE5s_etc to %struct.ETC addrspace(4)*))
 
   // anchor point separating constructors and destructors
   df(); // CHECK: call void @_Z2dfv()
 
   // Verify that we only call non-empty destructors
-  // CHECK-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %t_fa_ned)
-  // CHECK-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %t_f_ned)
-  // CHECK-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %t_b_ned)
-  // CHECK-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %vd)
-  // CHECK-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %ned)
-  // CHECK-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %ud)
-  // CHECK-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %ecd)
-  // CHECK-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %ed)
+  // NVPTX-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED* %t_fa_ned)
+  // NVPTX-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED* %t_f_ned)
+  // NVPTX-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED* %t_b_ned)
+  // NVPTX-NEXT: call void @_ZN2VDD1Ev(%struct.VD* %vd)
+  // NVPTX-NEXT: call void @_ZN3NEDD1Ev(%struct.NED* %ned)
+  // NVPTX-NEXT: call void @_ZN2UDD1Ev(%struct.UD* %ud)
+  // NVPTX-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD* %ecd)
+  // NVPTX-NEXT: call void @_ZN2EDD1Ev(%struct.ED* %ed)
+
+  // AMDGCN-NEXT: call void @_ZN8T_FA_NEDD1Ev(%struct.T_FA_NED addrspace(4)* %t_fa_ned)
+  // AMDGCN-NEXT: call void @_ZN7T_F_NEDD1Ev(%struct.T_F_NED addrspace(4)* %t_f_ned)
+  // AMDGCN-NEXT: call void @_ZN7T_B_NEDD1Ev(%struct.T_B_NED addrspace(4)* %t_b_ned)
+  // AMDGCN-NEXT: call void @_ZN2VDD1Ev(%struct.VD addrspace(4)* %vd)
+  // AMDGCN-NEXT: call void @_ZN3NEDD1Ev(%struct.NED addrspace(4)* %ned)
+  // AMDGCN-NEXT: call void @_ZN2UDD1Ev(%struct.UD addrspace(4)* %ud)
+  // AMDGCN-NEXT: call void @_ZN3ECDD1Ev(%struct.ECD addrspace(4)* %ecd)
+  // AMDGCN-NEXT: call void @_ZN2EDD1Ev(%struct.ED addrspace(4)* %ed)
 
   // CHECK-NEXT: ret void
 }
 
 // We should not emit global init function.
 // CHECK-NOT: @_GLOBAL__sub_I

test/CodeGenCUDA/device-vtable.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
 // Make sure we don't emit vtables for classes with methods that have
 // inappropriate target attributes. Currently it's mostly needed in
 // order to avoid emitting vtables for host-only classes on device
 // side where we can't codegen them.
 
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \
 // RUN:     | FileCheck %s -check-prefix=CHECK-HOST -check-prefix=CHECK-BOTH
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
 // RUN:     | FileCheck %s -check-prefix=CHECK-DEVICE -check-prefix=CHECK-BOTH
+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm -o - %s \
+// RUN:     | FileCheck %s -check-prefix=CHECK-DEVICE -check-prefix=CHECK-BOTH
 
 #include "Inputs/cuda.h"
 
 struct H  {
   virtual void method();
 };
 //CHECK-HOST: @_ZTV1H =
 //CHECK-HOST-SAME: @_ZN1H6methodEv

test/CodeGenCUDA/filter-decl.cu

 // RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - | FileCheck -check-prefix=CHECK-HOST %s
-// RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - -fcuda-is-device | FileCheck -check-prefix=CHECK-DEVICE %s
+// RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm %s -o - -fcuda-is-device | FileCheck -check-prefixes=CHECK-DEVICE,ITANIUM %s
+// RUN: %clang_cc1 -triple amdgcn -emit-llvm %s -o - -fcuda-is-device | FileCheck -check-prefixes=CHECK-DEVICE,AMDGCN %s
 
 #include "Inputs/cuda.h"
 
 // This has to be at the top of the file as that's where file-scope
 // asm ends up.
 // CHECK-HOST: module asm "file scope asm is host only"
 // CHECK-DEVICE-NOT: module asm "file scope asm is host only"
 __asm__("file scope asm is host only");
 
 // CHECK-HOST: constantdata = internal global
-// CHECK-DEVICE: constantdata = externally_initialized global
+// ITANIUM: constantdata = externally_initialized global
+// AMDGCN: constantdata = addrspace(2) externally_initialized global
 __constant__ char constantdata[256];
 
 // CHECK-HOST: devicedata = internal global
-// CHECK-DEVICE: devicedata = externally_initialized global
+// ITANIUM: devicedata = externally_initialized global
+// AMDGCN: devicedata = addrspace(1) externally_initialized global
 __device__ char devicedata[256];
 
 // CHECK-HOST: shareddata = internal global
-// CHECK-DEVICE: shareddata = global
+// ITANIUM: shareddata = global
+// AMDGCN: shareddata = addrspace(3) global
 __shared__ char shareddata[256];
 
 // CHECK-HOST: hostdata = global
 // CHECK-DEVICE-NOT: hostdata = global
 char hostdata[256];
 
 // CHECK-HOST: define{{.*}}implicithostonlyfunc
 // CHECK-DEVICE-NOT: define{{.*}}implicithostonlyfunc

test/CodeGenCUDA/function-overload.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
 // Make sure we handle target overloads correctly.  Most of this is checked in
 // sema, but special functions like constructors and destructors are here.
 //
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm -o - %s \
+// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
 
 #include "Inputs/cuda.h"
 
 // Check constructors/destructors for D/H functions
 int x;
 struct s_cd_dh {
   __host__ s_cd_dh() { x = 11; }
   __device__ s_cd_dh() { x = 12; }

test/CodeGenCUDA/kernel-args-alignment.cu

 // RUN: %clang_cc1 --std=c++11 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s | \
-// RUN:  FileCheck -check-prefix HOST -check-prefix CHECK %s
+// RUN:  FileCheck -check-prefixes=HOST,CHECK %s
 
 // RUN: %clang_cc1 --std=c++11 -fcuda-is-device -triple nvptx64-nvidia-cuda \
-// RUN:   -emit-llvm -o - %s | FileCheck -check-prefix DEVICE -check-prefix CHECK %s
+// RUN:   -emit-llvm -o - %s | FileCheck -check-prefixes=DEVICE,CHECK,NVPTX %s
+
+// RUN: %clang_cc1 --std=c++11 -fcuda-is-device -triple amdgcn-amd-amdhsa \
+// RUN:   -emit-llvm -o - %s -DAMDGCN| FileCheck -check-prefixes=DEVICE,CHECK,AMDGCN %s
 
 #include "Inputs/cuda.h"
 
 struct U {
   short x;
 } __attribute__((packed));
 
 struct S {
   int *ptr;
   char a;
   U u;
 };
 
 // Clang should generate a packed LLVM struct for S (denoted by the <>s),
 // otherwise this test isn't interesting.
-// CHECK: %struct.S = type <{ i32*, i8, %struct.U, [5 x i8] }>
+// HOST: %struct.S = type <{ i32*, i8, %struct.U, [5 x i8] }>
+// NVPTX: %struct.S = type <{ i32*, i8, %struct.U, [5 x i8] }>
+// AMDGCN: %struct.S = type <{ i32 addrspace(4)*, i8, %struct.U, [5 x i8] }>
 
 static_assert(alignof(S) == 8, "Unexpected alignment.");
 
 // HOST-LABEL: @_Z6kernelc1SPi
 // Marshalled kernel args should be:
 //   1. offset 0, width 1
 //   2. offset 8 (because alignof(S) == 8), width 16
 //   3. offset 24, width 8
 // HOST: call i32 @cudaSetupArgument({{[^,]*}}, i64 1, i64 0)
 // HOST: call i32 @cudaSetupArgument({{[^,]*}}, i64 16, i64 8)
 // HOST: call i32 @cudaSetupArgument({{[^,]*}}, i64 8, i64 24)
 
 // DEVICE-LABEL: @_Z6kernelc1SPi
-// DEVICE-SAME: i8{{[^,]*}}, %struct.S* byval align 8{{[^,]*}}, i32*
+// NVPTX-SAME: i8{{[^,]*}}, %struct.S* byval align 8{{[^,]*}}, i32*
+// AMDGCN-SAME: i8{{[^,]*}}, %struct.S* byval align 8{{[^,]*}}, i32 addrspace(4)*
 __global__ void kernel(char a, S s, int *b) {}

test/CodeGenCUDA/llvm-used.cu

-// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx64-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx64-unknown-unknown | FileCheck -check-prefix=NVPTX %s
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple amdgcn-amd-amdhsa | FileCheck -check-prefix=AMDGCN %s
 
 
 // Make sure we emit the proper addrspacecast for llvm.used.  PR22383 exposed an
 // issue where we were generating a bitcast instead of an addrspacecast.
 
-// CHECK: @llvm.used = appending global [1 x i8*] [i8* addrspacecast (i8 addrspace(1)* bitcast ([0 x i32] addrspace(1)* @a to i8 addrspace(1)*) to i8*)], section "llvm.metadata"
+// NVPTX: @llvm.used = appending global [1 x i8*] [i8* addrspacecast (i8 addrspace(1)* bitcast ([0 x i32] addrspace(1)* @a to i8 addrspace(1)*) to i8*)], section "llvm.metadata"
+// AMDGCN: @llvm.used = appending global [1 x i8 addrspace(4)*] [i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast ([0 x i32] addrspace(1)* @a to i8 addrspace(1)*) to i8 addrspace(4)*)], section "llvm.metadata"
 __attribute__((device)) __attribute__((__used__)) int a[] = {};

test/CodeGenCUDA/printf.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm \
-// RUN:   -o - %s | FileCheck %s
+// RUN:   -o - %s | FileCheck -check-prefixes=CHECK,NVPTX %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm \
+// RUN:   -o - %s | FileCheck -check-prefixes=CHECK,AMDGCN %s
 
 #include "Inputs/cuda.h"
 
 extern "C" __device__ int vprintf(const char*, const char*);
 
 // Check a simple call to printf end-to-end.
 // CHECK: [[SIMPLE_PRINTF_TY:%[a-zA-Z0-9_]+]] = type { i32, i64, double }
+// CHECK-LABEL: define i32 @_Z11CheckSimplev()
 __device__ int CheckSimple() {
-  // CHECK: [[BUF:%[a-zA-Z0-9_]+]] = alloca [[SIMPLE_PRINTF_TY]]
+  // NVPTX: [[BUF:%[a-zA-Z0-9_]+]] = alloca [[SIMPLE_PRINTF_TY]]
+  // AMDGCN: [[ALLOCA:%[a-zA-Z0-9_]+]] = alloca [[SIMPLE_PRINTF_TY]]
+  // AMDGCN: [[BUF:%[a-zA-Z0-9_]+]] = addrspacecast %printf_args* [[ALLOCA]] to %printf_args addrspace(4)*
   // CHECK: [[FMT:%[0-9]+]] = load{{.*}}%fmt
   const char* fmt = "%d %lld %f";
-  // CHECK: [[PTR0:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]* [[BUF]], i32 0, i32 0
-  // CHECK: store i32 1, i32* [[PTR0]], align 4
-  // CHECK: [[PTR1:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]* [[BUF]], i32 0, i32 1
-  // CHECK: store i64 2, i64* [[PTR1]], align 8
-  // CHECK: [[PTR2:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]* [[BUF]], i32 0, i32 2
-  // CHECK: store double 3.0{{[^,]*}}, double* [[PTR2]], align 8
-  // CHECK: [[BUF_CAST:%[0-9]+]] = bitcast [[SIMPLE_PRINTF_TY]]* [[BUF]] to i8*
-  // CHECK: [[RET:%[0-9]+]] = call i32 @vprintf(i8* [[FMT]], i8* [[BUF_CAST]])
+  // CHECK: [[PTR0:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]{{.*}}* [[BUF]], i32 0, i32 0
+  // CHECK: store i32 1, i32{{.*}}* [[PTR0]], align 4
+  // CHECK: [[PTR1:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]{{.*}}* [[BUF]], i32 0, i32 1
+  // CHECK: store i64 2, i64{{.*}}* [[PTR1]], align 8
+  // CHECK: [[PTR2:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]{{.*}}* [[BUF]], i32 0, i32 2
+  // CHECK: store double 3.0{{[^,]*}}, double{{.*}}* [[PTR2]], align 8
+  // CHECK: [[BUF_CAST:%[0-9]+]] = bitcast [[SIMPLE_PRINTF_TY]]{{.*}}* [[BUF]] to i8{{.*}}*
+  // CHECK: [[RET:%[0-9]+]] = call i32 @vprintf(i8{{.*}}* [[FMT]], i8{{.*}}* [[BUF_CAST]])
   // CHECK: ret i32 [[RET]]
   return printf(fmt, 1, 2ll, 3.0);
 }
 
+// CHECK-LABEL: define void @_Z11CheckNoArgsv()
 __device__ void CheckNoArgs() {
-  // CHECK: call i32 @vprintf({{.*}}, i8* null){{$}}
+  // CHECK: call i32 @vprintf({{.*}}, i8{{.*}}* null){{$}}
   printf("hello, world!");
 }
 
 // Check that printf's alloca happens in the entry block, not inside the if
 // statement.
 __device__ bool foo();
+// CHECK-LABEL: define void @_Z25CheckAllocaIsInEntryBlockv()
 __device__ void CheckAllocaIsInEntryBlock() {
   // CHECK: alloca %printf_args
   // CHECK: call {{.*}} @_Z3foov()
   if (foo()) {
     printf("%d", 42);
   }
 }

test/CodeGenCXX/amdgcn-global-init.cpp

+// RUN: %clang_cc1 -std=c++11 -triple=amdgcn-amd-amdhsa -emit-llvm -fexceptions %s -o - |FileCheck %s
+// RUN: %clang_cc1 -std=c++11 -triple=amdgcn-amd-amdhsa -emit-llvm %s -o - |FileCheck -check-prefix CHECK-NOEXC %s
+// RUN: %clang_cc1 -std=c++11 -triple=amdgcn-amd-amdhsa -emit-llvm \
+// RUN:     -momit-leaf-frame-pointer -mdisable-fp-elim %s -o - \
+// RUN:   | FileCheck -check-prefix CHECK-FP %s
+
+struct A {
+  A();
+  ~A();
+};
+
+struct B { B(); ~B(); };
+
+struct C { void *field; };
+
+struct D { ~D(); };
+
+// CHECK: @__dso_handle = external hidden addrspace(1) global i8
+// CHECK: @c = addrspace(1) global %struct.C zeroinitializer, align 8
+
+// PR6205: The casts should not require global initializers
+// CHECK: @_ZN6PR59741cE = external addrspace(1) global %"struct.PR5974::C"
+// CHECK: @_ZN6PR59741aE = addrspace(1) global %"struct.PR5974::A" addrspace(4)* addrspacecast (%"struct.PR5974::A" addrspace(1)* getelementptr inbounds (%"struct.PR5974::C", %"struct.PR5974::C" addrspace(1)* @_ZN6PR59741cE, i32 0, i32 0) to %"struct.PR5974::A" addrspace(4)*), align 8
+// CHECK: @_ZN6PR59741bE = addrspace(1) global %"struct.PR5974::B" addrspace(4)* bitcast (i8 addrspace(4)* getelementptr (i8, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"struct.PR5974::C" addrspace(1)* @_ZN6PR59741cE to i8 addrspace(1)*) to i8 addrspace(4)*), i64 4) to %"struct.PR5974::B" addrspace(4)*), align 8
+
+// CHECK: call void @_ZN1AC1Ev(%struct.A addrspace(4)* addrspacecast (%struct.A addrspace(1)* @a to %struct.A addrspace(4)*))
+// CHECK: call i32 @__cxa_atexit(void (i8 addrspace(4)*)* bitcast (void (%struct.A addrspace(4)*)* @_ZN1AD1Ev to void (i8 addrspace(4)*)*), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds (%struct.A, %struct.A addrspace(1)* @a, i32 0, i32 0) to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* @__dso_handle to i8 addrspace(4)*))
+A a;
+
+// CHECK: call void @_ZN1BC1Ev(%struct.B addrspace(4)* addrspacecast (%struct.B addrspace(1)* @b to %struct.B addrspace(4)*))
+// CHECK: call i32 @__cxa_atexit(void (i8 addrspace(4)*)* bitcast (void (%struct.B addrspace(4)*)* @_ZN1BD1Ev to void (i8 addrspace(4)*)*), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds (%struct.B, %struct.B addrspace(1)* @b, i32 0, i32 0) to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* @__dso_handle to i8 addrspace(4)*))
+B b;
+
+// PR6205: this should not require a global initializer
+// CHECK-NOT: call void @_ZN1CC1Ev
+C c;
+
+// CHECK: call i32 @__cxa_atexit(void (i8 addrspace(4)*)* bitcast (void (%struct.D addrspace(4)*)* @_ZN1DD1Ev to void (i8 addrspace(4)*)*), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* getelementptr inbounds (%struct.D, %struct.D addrspace(1)* @d, i32 0, i32 0) to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* @__dso_handle to i8 addrspace(4)*))
+D d;
+
+// <rdar://problem/7458115>
+namespace test1 {
+  int f();
+  const int x = f();   // This has side-effects and gets emitted immediately.
+  const int y = x - 1; // This gets deferred.
+  const int z = ~y;    // This also gets deferred, but gets "undeferred" before y.
+  int test() { return z; }
+// CHECK-LABEL:      define i32 @_ZN5test14testEv()
+
+  // All of these initializers end up delayed, so we check them later.
+}
+
+// <rdar://problem/8246444>
+namespace test2 {
+  struct allocator { allocator(); ~allocator(); };
+  struct A { A(const allocator &a = allocator()); ~A(); };
+
+  A a;
+// CHECK: call void @_ZN5test29allocatorC1Ev(
+// CHECK: invoke void @_ZN5test21AC1ERKNS_9allocatorE(
+// CHECK: call void @_ZN5test29allocatorD1Ev(
+// CHECK: call i32 @__cxa_atexit({{.*}} @_ZN5test21AD1Ev {{.*}} @_ZN5test21aE
+}
+
+namespace test3 {
+  // Tested at the beginning of the file.
+  const char * const var = "string";
+  extern const char * const var;
+
+  const char *test() { return var; }
+}
+
+namespace test4 {
+  struct A {
+    A();
+  };
+  extern int foo();
+
+  // This needs an initialization function and guard variables.
+  // CHECK: load i8, i8 addrspace(1)* bitcast (i64 addrspace(1)* @_ZGVN5test41xE to i8 addrspace(1)*)
+  // CHECK: [[CALL:%.*]] = call i32 @_ZN5test43fooEv
+  // CHECK-NEXT: store i32 %call, i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @_ZN5test41xE to i32 addrspace(4)*)
+  // CHECK-NEXT: store i64 1, i64 addrspace(1)* @_ZGVN5test41xE
+  __attribute__((weak)) int x = foo();
+}
+
+namespace PR5974 {
+  struct A { int a; };
+  struct B { int b; };
+  struct C : A, B { int c; };
+
+  extern C c;
+
+  // These should not require global initializers.
+  A* a = &c;
+  B* b = &c;
+}
+
+// PR9570: the indirect field shouldn't crash IR gen.
+namespace test5 {
+  static union {
+    unsigned bar[4096] __attribute__((aligned(128)));
+  };
+}
+
+namespace std { struct type_info; }
+
+namespace test6 {
+  struct A { virtual ~A(); };
+  struct B : A {};
+  extern A *p;
+
+  // We must emit a dynamic initializer for 'q', because it could throw.
+  B *const q = &dynamic_cast<B&>(*p);
+  // CHECK: call void @__cxa_bad_cast()
+  // CHECK: store {{.*}} @_ZN5test6L1qE
+
+  // We don't need to emit 'r' at all, because it has internal linkage, is
+  // unused, and its initialization has no side-effects.
+  B *const r = dynamic_cast<B*>(p);
+  // CHECK-NOT: call void @__cxa_bad_cast()
+  // CHECK-NOT: store {{.*}} @_ZN5test6L1rE
+
+  // This can throw, so we need to emit it.
+  const std::type_info *const s = &typeid(*p);
+  // CHECK: store {{.*}} @_ZN5test6L1sE
+
+  // This can't throw, so we don't.
+  const std::type_info *const t = &typeid(p);
+  // CHECK-NOT: @_ZN5test6L1tE
+
+  extern B *volatile v;
+  // CHECK: store {{.*}} @_ZN5test6L1wE
+  B *const w = dynamic_cast<B*>(v);
+
+  // CHECK: load volatile
+  // CHECK: store {{.*}} @_ZN5test6L1xE
+  const int x = *(volatile int*)0x1234;
+
+  namespace {
+    int a = int();
+    volatile int b = int();
+    int c = a;
+    int d = b;
+    // CHECK-NOT: store {{.*}} @_ZN5test6{{[A-Za-z0-9_]*}}1aE
+    // CHECK-NOT: store {{.*}} @_ZN5test6{{[A-Za-z0-9_]*}}1bE
+    // CHECK-NOT: store {{.*}} @_ZN5test6{{[A-Za-z0-9_]*}}1cE
+    // CHECK: load volatile {{.*}} @_ZN5test6{{[A-Za-z0-9_]*}}1bE
+    // CHECK: store {{.*}} @_ZN5test6{{[A-Za-z0-9_]*}}1dE
+  }
+}
+
+namespace test7 {
+  struct A { A(); };
+  struct B { ~B(); int n; };
+  struct C { C() = default; C(const C&); int n; };
+  struct D {};
+
+  // CHECK: call void @_ZN5test71AC1Ev({{.*}}@_ZN5test7L1aE
+  const A a = A();
+
+  // CHECK: call i32 @__cxa_atexit({{.*}} @_ZN5test71BD1Ev{{.*}} @_ZN5test7L2b1E
+  // CHECK: call i32 @__cxa_atexit({{.*}} @_ZN5test71BD1Ev{{.*}} @_ZGRN5test72b2E
+  // CHECK: call void @_ZN5test71BD1Ev(
+  // CHECK: store {{.*}} @_ZN5test7L2b3E
+  const B b1 = B();
+  const B &b2 = B();
+  const int b3 = B().n;
+
+  // CHECK-NOT: @_ZN5test7L2c1E
+  // CHECK: call void @llvm.memset{{.*}} @_ZN5test7L2c1E
+  // CHECK-NOT: @_ZN5test7L2c1E
+  // CHECK: @_ZN5test7L2c2E
+  // CHECK-NOT: @_ZN5test7L2c3E
+  // CHECK: @_ZN5test7L2c4E
+  const C c1 = C();
+  const C c2 = static_cast<const C&>(C());
+  const int c3 = C().n;
+  const int c4 = C(C()).n;
+
+  // CHECK-NOT: @_ZN5test7L1dE
+  const D d = D();
+
+  // CHECK: store {{.*}} @_ZN5test71eE
+  int f(), e = f();
+}
+
+
+// At the end of the file, we check that y is initialized before z.
+
+// CHECK:      define internal void [[TEST1_Z_INIT:@.*]]()
+// CHECK:        load i32, i32 addrspace(4)* addrspacecast (i32 addrspace(2)* @_ZN5test1L1yE to i32 addrspace(4)*)
+// CHECK-NEXT:   xor
+// CHECK-NEXT:   store i32 {{.*}}, i32 addrspace(4)* addrspacecast (i32 addrspace(2)* @_ZN5test1L1zE to i32 addrspace(4)*)
+// CHECK:      define internal void [[TEST1_Y_INIT:@.*]]()
+// CHECK:        load i32, i32 addrspace(4)* addrspacecast (i32 addrspace(2)* @_ZN5test1L1xE to i32 addrspace(4)*)
+// CHECK-NEXT:   sub
+// CHECK-NEXT:   store i32 {{.*}}, i32 addrspace(4)* addrspacecast (i32 addrspace(2)* @_ZN5test1L1yE to i32 addrspace(4)*)
+
+// CHECK: define internal void @_GLOBAL__sub_I_amdgcn_global_init.cpp() #{{[0-9]+}}
+// CHECK:   call void [[TEST1_Y_INIT]]
+// CHECK:   call void [[TEST1_Z_INIT]]
+
+// rdar://problem/8090834: this should be nounwind
+// CHECK-NOEXC: define internal void @_GLOBAL__sub_I_amdgcn_global_init.cpp() [[NUW:#[0-9]+]]
+
+// CHECK-NOEXC: attributes [[NUW]] = { noinline nounwind{{.*}} }
+
+// PR21811: attach the appropriate attribute to the global init function
+// CHECK-FP: define internal void @_GLOBAL__sub_I_amdgcn_global_init.cpp() [[NUX:#[0-9]+]]
+// CHECK-FP: attributes [[NUX]] = { noinline nounwind {{.*}}"no-frame-pointer-elim-non-leaf"{{.*}} }

test/OpenMP/nvptx_parallel_codegen.cpp

 // Test target codegen - host bc file has to be created first.
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=amdgcn -emit-llvm-bc %s -o %t-x86-host.bc
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
 // RUN: %clang_cc1 -verify -fopenmp -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 
 template<typename tx>
 tx ftemplate(int n) {

 int bar(int n){
   int a = 0;
 
   a += ftemplate<int>(n);
 
   return a;
 }
 
-  // CHECK-NOT: define {{.*}}void {{@__omp_offloading_.+template.+l17}}_worker()
+  // CHECK-NOT: define {{.*}}void {{@__omp_offloading_.+template.+l18}}_worker()
 
 
 
 
 
 
-  // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l26}}_worker()
+  // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l27}}_worker()
   // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
   // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
   // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
   // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
   // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
   //
   // CHECK: [[AWAIT_WORK]]
   // CHECK: call void @llvm.nvvm.barrier0()
int bar(int n){
   //
   // CHECK: [[BAR_PARALLEL]]
   // CHECK: call void @llvm.nvvm.barrier0()
   // CHECK: br label {{%?}}[[AWAIT_WORK]]
   //
   // CHECK: [[EXIT]]
   // CHECK: ret void
 
-  // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l26]](i[[SZ:32|64]]
+  // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l27]](i[[SZ:32|64]]
   // Create local storage for each capture.
   // CHECK:  [[LOCAL_A:%.+]] = alloca i[[SZ]],
   // CHECK-DAG:  store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
   // Store captures in the context.
   // CHECK-64-DAG:[[REF_A:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
   //
   // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
   // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
int bar(int n){
   // CHECK: ret void
 
 
 
 
 
 
 
-  // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l43}}_worker()
+  // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l44}}_worker()
   // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
   // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
   // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
   // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
   // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
   //
   // CHECK: [[AWAIT_WORK]]
   // CHECK: call void @llvm.nvvm.barrier0()
int bar(int n){
   //
   // CHECK: [[BAR_PARALLEL]]
   // CHECK: call void @llvm.nvvm.barrier0()
   // CHECK: br label {{%?}}[[AWAIT_WORK]]
   //
   // CHECK: [[EXIT]]
   // CHECK: ret void
 
-  // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l43]](i[[SZ:32|64]]
+  // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l44]](i[[SZ:32|64]]
   // Create local storage for each capture.
   // CHECK:  [[LOCAL_N:%.+]] = alloca i[[SZ]],
   // CHECK:  [[LOCAL_A:%.+]] = alloca i[[SZ]],
   // CHECK:  [[LOCAL_AA:%.+]] = alloca i[[SZ]],
   // CHECK:  [[LOCAL_B:%.+]] = alloca [10 x i32]*
   // CHECK-DAG:  store i[[SZ]] [[ARG_N:%.+]], i[[SZ]]* [[LOCAL_N]]
   // CHECK-DAG:  store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
   // CHECK-DAG:  store i[[SZ]] [[ARG_AA:%.+]], i[[SZ]]* [[LOCAL_AA]]