[AMDGPU] Fix size and alignment of size_t and pointer types

[yaxunl]

Summary

This patch is split from https://reviews.llvm.org/D33842

[rjmccall]

Looks good.

lib/Basic/Targets.cpp

                     : DataLayoutStringR600);
     assert(DataLayout->getAllocaAddrSpace() == AS.Private);
 
     setAddressSpaceMap(Triple.getOS() == llvm::Triple::Mesa3D ||
                        Triple.getEnvironment() == llvm::Triple::OpenCL ||
                        Triple.getEnvironmentName() == "amdgizcl" ||
                        !isAMDGCN(Triple));
     UseAddrSpaceMapMangling = true;
+
+    // Set pointer width and alignment for target address space 0.
+    PointerWidth = PointerAlign = DataLayout->getPointerSizeInBits();
+    if (getMaxPointerWidth() == 64) {
+      LongWidth = LongAlign = 64;
+      SizeType = UnsignedLong;
+      PtrDiffType = SignedLong;
+      IntPtrType = SignedLong;
+    }
   }
 
   void setAddressSpaceMap(bool DefaultIsPrivate) {
     if (isGenericZero(getTriple())) {
       AddrSpaceMap = DefaultIsPrivate ? &AMDGPUGenIsZeroDefIsPrivMap
                                       : &AMDGPUGenIsZeroDefIsGenMap;
     } else {
       AddrSpaceMap = DefaultIsPrivate ? &AMDGPUPrivIsZeroDefIsPrivMap

       return 32;
 
     if (AddrSpace == AS.Private || AddrSpace == AS.Local) {
       return 32;
     }
     return 64;
   }
 
+  uint64_t getPointerAlignV(unsigned AddrSpace) const override {
+    return getPointerWidthV(AddrSpace);
+  }
+
   uint64_t getMaxPointerWidth() const override {
     return getTriple().getArch() == llvm::Triple::amdgcn ? 64 : 32;
   }
 
   const char * getClobbers() const override {
     return "";
   }
 

     }
   }
 
   LangAS::ID getOpenCLImageAddrSpace() const override {
     return LangAS::opencl_constant;
   }
 
   /// \returns Target specific vtbl ptr address space.
-  unsigned getVtblPtrAddressSpace() const override {
-    // \todo: We currently have address spaces defined in AMDGPU Backend. It
-    // would be nice if we could use it here instead of using bare numbers (same
-    // applies to getDWARFAddressSpace).
-    return 2; // constant.
-  }
+  unsigned getVtblPtrAddressSpace() const override { return AS.Constant; }
 
   /// \returns If a target requires an address within a target specific address
   /// space \p AddressSpace to be converted in order to be used, then return the
   /// corresponding target specific DWARF address space.
   ///
   /// \returns Otherwise return None and no conversion will be emitted in the
   /// DWARF.
   Optional<unsigned> getDWARFAddressSpace(

test/CodeGen/default-address-space.c

 // PIZ-LABEL: define void @test4(i32 addrspace(4)* %a)
 // PIZ: %[[alloca:.*]] = alloca i32 addrspace(4)*
 // PIZ: %[[a_addr:.*]] = addrspacecast{{.*}} %[[alloca]] to i32 addrspace(4)* addrspace(4)*
 // PIZ: store i32 addrspace(4)* %a, i32 addrspace(4)* addrspace(4)* %[[a_addr]]
 // PIZ: %[[r0:.*]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* %[[a_addr]]
 // PIZ: %[[arrayidx:.*]] = getelementptr inbounds i32, i32 addrspace(4)* %[[r0]]
 // PIZ: store i32 0, i32 addrspace(4)* %[[arrayidx]]
 // CHECK-LABEL: define void @test4(i32* %a)
-// CHECK: %[[alloca:.*]] = alloca i32*, align 4, addrspace(5)
+// CHECK: %[[alloca:.*]] = alloca i32*, align 8, addrspace(5)
 // CHECK: %[[a_addr:.*]] = addrspacecast{{.*}} %[[alloca]] to i32**
 // CHECK: store i32* %a, i32** %[[a_addr]]
 // CHECK: %[[r0:.*]] = load i32*, i32** %[[a_addr]]
 // CHECK: %[[arrayidx:.*]] = getelementptr inbounds i32, i32* %[[r0]]
 // CHECK: store i32 0, i32* %[[arrayidx]]
 void test4(int *a) {
   a[0] = 0;
 }

test/CodeGenCXX/amdgcn-automatic-variable.cpp

void func1(int *x) {
   *x = 1;
 }
 
 // CHECK-LABEL: define void @_Z5func2v()
 void func2(void) {
   // CHECK: %lv1 = alloca i32, align 4, addrspace(5)
   // CHECK: %lv2 = alloca i32, align 4, addrspace(5)
   // CHECK: %la = alloca [100 x i32], align 4, addrspace(5)
-  // CHECK: %lp1 = alloca i32*, align 4, addrspace(5)
-  // CHECK: %lp2 = alloca i32*, align 4, addrspace(5)
+  // CHECK: %lp1 = alloca i32*, align 8, addrspace(5)
+  // CHECK: %lp2 = alloca i32*, align 8, addrspace(5)
   // CHECK: %lvc = alloca i32, align 4, addrspace(5)
 
   // CHECK: %[[r0:.*]] = addrspacecast i32 addrspace(5)* %lv1 to i32*
   // CHECK: store i32 1, i32* %[[r0]]
   int lv1;
   lv1 = 1;
   // CHECK: %[[r1:.*]] = addrspacecast i32 addrspace(5)* %lv2 to i32*
   // CHECK: store i32 2, i32* %[[r1]]
   int lv2 = 2;
 
   // CHECK: %[[r2:.*]] = addrspacecast [100 x i32] addrspace(5)* %la to [100 x i32]*
   // CHECK: %[[arrayidx:.*]] = getelementptr inbounds [100 x i32], [100 x i32]* %[[r2]], i64 0, i64 0
   // CHECK: store i32 3, i32* %[[arrayidx]], align 4
   int la[100];
   la[0] = 3;
 
   // CHECK: %[[r3:.*]] = addrspacecast i32* addrspace(5)* %lp1 to i32**
-  // CHECK: store i32* %[[r0]], i32** %[[r3]], align 4
+  // CHECK: store i32* %[[r0]], i32** %[[r3]], align 8
   int *lp1 = &lv1;
 
   // CHECK: %[[r4:.*]] = addrspacecast i32* addrspace(5)* %lp2 to i32**
   // CHECK: %[[arraydecay:.*]] = getelementptr inbounds [100 x i32], [100 x i32]* %[[r2]], i32 0, i32 0
-  // CHECK: store i32* %[[arraydecay]], i32** %[[r4]], align 4
+  // CHECK: store i32* %[[arraydecay]], i32** %[[r4]], align 8
   int *lp2 = la;
 
   // CHECK: call void @_Z5func1Pi(i32* %[[r0]])
   func1(&lv1);
 
   // CHECK: %[[r5:.*]] = addrspacecast i32 addrspace(5)* %lvc to i32*
   // CHECK: store i32 4, i32* %[[r5]]
   // CHECK: store i32 4, i32* %[[r0]]

 // CHECK-LABEL: define void @_Z5func4i
 void func4(int x) {
   // CHECK: %[[x_addr:.*]] = alloca i32, align 4, addrspace(5)
   // CHECK: %[[r0:.*]] = addrspacecast i32 addrspace(5)* %[[x_addr]] to i32*
   // CHECK: store i32 %x, i32* %[[r0]], align 4
   // CHECK: call void @_Z5func1Pi(i32* %[[r0]])
   func1(&x);
 }
+
+// CHECK-NOT: !opencl.ocl.version

test/CodeGenOpenCL/amdgcn-automatic-variable.cl

 
 // CHECK-LABEL: define void @func2()
 void func2(void) {
   // CHECK: %lv1 = alloca i32, align 4, addrspace(5)
   // CHECK: %lv2 = alloca i32, align 4, addrspace(5)
   // CHECK: %la = alloca [100 x i32], align 4, addrspace(5)
   // CL12: %lp1 = alloca i32 addrspace(5)*, align 4, addrspace(5)
   // CL12: %lp2 = alloca i32 addrspace(5)*, align 4, addrspace(5)
-  // CL20: %lp1 = alloca i32*, align 4, addrspace(5)
-  // CL20: %lp2 = alloca i32*, align 4, addrspace(5)
+  // CL20: %lp1 = alloca i32*, align 8, addrspace(5)
+  // CL20: %lp2 = alloca i32*, align 8, addrspace(5)
   // CHECK: %lvc = alloca i32, align 4, addrspace(5)
 
   // CHECK: store i32 1, i32 addrspace(5)* %lv1
   int lv1;
   lv1 = 1;
   // CHECK: store i32 2, i32 addrspace(5)* %lv2
   int lv2 = 2;
 
   // CHECK: %[[arrayidx:.*]] = getelementptr inbounds [100 x i32], [100 x i32] addrspace(5)* %la, i64 0, i64 0
   // CHECK: store i32 3, i32 addrspace(5)* %[[arrayidx]], align 4
   int la[100];
   la[0] = 3;
 
   // CL12: store i32 addrspace(5)* %lv1, i32 addrspace(5)* addrspace(5)* %lp1, align 4
   // CL20: %[[r0:.*]] = addrspacecast i32 addrspace(5)* %lv1 to i32*
-  // CL20: store i32* %[[r0]], i32* addrspace(5)* %lp1, align 4
+  // CL20: store i32* %[[r0]], i32* addrspace(5)* %lp1, align 8
   int *lp1 = &lv1;
 
   // CHECK: %[[arraydecay:.*]] = getelementptr inbounds [100 x i32], [100 x i32] addrspace(5)* %la, i32 0, i32 0
   // CL12: store i32 addrspace(5)* %[[arraydecay]], i32 addrspace(5)* addrspace(5)* %lp2, align 4
   // CL20: %[[r1:.*]] = addrspacecast i32 addrspace(5)* %[[arraydecay]] to i32*
-  // CL20: store i32* %[[r1]], i32* addrspace(5)* %lp2, align 4
+  // CL20: store i32* %[[r1]], i32* addrspace(5)* %lp2, align 8
   int *lp2 = la;
 
   // CL12: call void @func1(i32 addrspace(5)* %lv1)
   // CL20: %[[r2:.*]] = addrspacecast i32 addrspace(5)* %lv1 to i32*
   // CL20: call void @func1(i32* %[[r2]])
   func1(&lv1);
 
   // CHECK: store i32 4, i32 addrspace(5)* %lvc
   // CHECK: store i32 4, i32 addrspace(5)* %lv1
   const int lvc = 4;
   lv1 = lvc;
 }

test/CodeGenOpenCL/amdgpu-nullptr.cl

 // Test 0 as initializer.
 
 // CHECK: @private_p = local_unnamed_addr addrspace(1) global i8* null, align 4
 private char *private_p = 0;
 
 // CHECK: @local_p = local_unnamed_addr addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4
 local char *local_p = 0;
 
-// CHECK: @global_p = local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 4
+// CHECK: @global_p = local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 8
 global char *global_p = 0;
 
-// CHECK: @constant_p = local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 4
+// CHECK: @constant_p = local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 8
 constant char *constant_p = 0;
 
-// CHECK: @generic_p = local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 4
+// CHECK: @generic_p = local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 8
 generic char *generic_p = 0;
 
 // Test NULL as initializer.
 
 // CHECK: @private_p_NULL = local_unnamed_addr addrspace(1) global i8* null, align 4
 private char *private_p_NULL = NULL;
 
 // CHECK: @local_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4
 local char *local_p_NULL = NULL;
 
-// CHECK: @global_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 4
+// CHECK: @global_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 8
 global char *global_p_NULL = NULL;
 
-// CHECK: @constant_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 4
+// CHECK: @constant_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 8
 constant char *constant_p_NULL = NULL;
 
-// CHECK: @generic_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 4
+// CHECK: @generic_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 8
 generic char *generic_p_NULL = NULL;
 
 // Test constant folding of null pointer.
 // A null pointer should be folded to a null pointer in the target address space.
 
-// CHECK: @fold_generic = local_unnamed_addr addrspace(1) global i32 addrspace(4)* null, align 4
+// CHECK: @fold_generic = local_unnamed_addr addrspace(1) global i32 addrspace(4)* null, align 8
 generic int *fold_generic = (global int*)(generic float*)(private char*)0;
 
 // CHECK: @fold_priv = local_unnamed_addr addrspace(1) global i16* null, align 4
 private short *fold_priv = (private short*)(generic int*)(global void*)0;
 
 // CHECK: @fold_priv_arith = local_unnamed_addr addrspace(1) global i8* inttoptr (i32 10 to i8*), align 4
 private char *fold_priv_arith = (private char*)0 + 10;
 

 
 // Test static variable initialization.
 
 // NOOPT: @test_static_var_private.sp1 = internal addrspace(1) global i8* null, align 4
 // NOOPT: @test_static_var_private.sp2 = internal addrspace(1) global i8* null, align 4
 // NOOPT: @test_static_var_private.sp3 = internal addrspace(1) global i8* null, align 4
 // NOOPT: @test_static_var_private.sp4 = internal addrspace(1) global i8* null, align 4
 // NOOPT: @test_static_var_private.sp5 = internal addrspace(1) global i8* null, align 4
-// NOOPT: @test_static_var_private.SS1 = internal addrspace(1) global %struct.StructTy1 { i8* null, i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }, align 4
-// NOOPT: @test_static_var_private.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 4
+// NOOPT: @test_static_var_private.SS1 = internal addrspace(1) global %struct.StructTy1 { i8* null, i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }, align 8
+// NOOPT: @test_static_var_private.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 8
 
 void test_static_var_private(void) {
   static private char *sp1 = 0;
   static private char *sp2 = NULL;
   static private char *sp3;
   static private char *sp4 = (private char*)((void)0, 0);
   const int x = 0;
   static private char *sp5 = (private char*)x;
   static StructTy1 SS1;
   static StructTy2 SS2;
 }
 
 // NOOPT: @test_static_var_local.sp1 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4
 // NOOPT: @test_static_var_local.sp2 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4
 // NOOPT: @test_static_var_local.sp3 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4
 // NOOPT: @test_static_var_local.sp4 = internal addrspace(1) global i8 addrspace(3)* null, align 4
 // NOOPT: @test_static_var_local.sp5 = internal addrspace(1) global i8 addrspace(3)* null, align 4
-// NOOPT: @test_static_var_local.SS1 = internal addrspace(1) global %struct.StructTy1 { i8* null, i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }, align 4
-// NOOPT: @test_static_var_local.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 4
+// NOOPT: @test_static_var_local.SS1 = internal addrspace(1) global %struct.StructTy1 { i8* null, i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }, align 8
+// NOOPT: @test_static_var_local.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 8
 void test_static_var_local(void) {
   static local char *sp1 = 0;
   static local char *sp2 = NULL;
   static local char *sp3;
   static local char *sp4 = (local char*)((void)0, 0);
   const int x = 0;
   static local char *sp5 = (local char*)x;
   static StructTy1 SS1;
   static StructTy2 SS2;
 }
 
 // Test function-scope variable initialization.
 // NOOPT-LABEL: @test_func_scope_var_private(
 // NOOPT: store i8* null, i8** %sp1, align 4
 // NOOPT: store i8* null, i8** %sp2, align 4
 // NOOPT: store i8* null, i8** %sp3, align 4
 // NOOPT: store i8* null, i8** %sp4, align 4
 // NOOPT: %[[SS1:.*]] = bitcast %struct.StructTy1* %SS1 to i8*
-// NOOPT: call void @llvm.memcpy.p0i8.p2i8.i64(i8* %[[SS1]], i8 addrspace(2)* bitcast (%struct.StructTy1 addrspace(2)* @test_func_scope_var_private.SS1 to i8 addrspace(2)*), i64 32, i32 4, i1 false)
+// NOOPT: call void @llvm.memcpy.p0i8.p2i8.i64(i8* %[[SS1]], i8 addrspace(2)* bitcast (%struct.StructTy1 addrspace(2)* @test_func_scope_var_private.SS1 to i8 addrspace(2)*), i64 32, i32 8, i1 false)
 // NOOPT: %[[SS2:.*]] = bitcast %struct.StructTy2* %SS2 to i8*
-// NOOPT: call void @llvm.memset.p0i8.i64(i8* %[[SS2]], i8 0, i64 24, i32 4, i1 false)
+// NOOPT: call void @llvm.memset.p0i8.i64(i8* %[[SS2]], i8 0, i64 24, i32 8, i1 false)
 void test_func_scope_var_private(void) {
   private char *sp1 = 0;
   private char *sp2 = NULL;
   private char *sp3 = (private char*)((void)0, 0);
   const int x = 0;
   private char *sp4 = (private char*)x;
   StructTy1 SS1 = {0, 0, 0, 0, 0};
   StructTy2 SS2 = {0, 0, 0};
 }
 
 // Test function-scope variable initialization.
 // NOOPT-LABEL: @test_func_scope_var_local(
 // NOOPT: store i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(3)** %sp1, align 4
 // NOOPT: store i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(3)** %sp2, align 4
 // NOOPT: store i8 addrspace(3)* null, i8 addrspace(3)** %sp3, align 4
 // NOOPT: store i8 addrspace(3)* null, i8 addrspace(3)** %sp4, align 4
 // NOOPT: %[[SS1:.*]] = bitcast %struct.StructTy1* %SS1 to i8*
-// NOOPT: call void @llvm.memcpy.p0i8.p2i8.i64(i8* %[[SS1]], i8 addrspace(2)* bitcast (%struct.StructTy1 addrspace(2)* @test_func_scope_var_local.SS1 to i8 addrspace(2)*), i64 32, i32 4, i1 false)
+// NOOPT: call void @llvm.memcpy.p0i8.p2i8.i64(i8* %[[SS1]], i8 addrspace(2)* bitcast (%struct.StructTy1 addrspace(2)* @test_func_scope_var_local.SS1 to i8 addrspace(2)*), i64 32, i32 8, i1 false)
 // NOOPT: %[[SS2:.*]] = bitcast %struct.StructTy2* %SS2 to i8*
-// NOOPT: call void @llvm.memset.p0i8.i64(i8* %[[SS2]], i8 0, i64 24, i32 4, i1 false)
+// NOOPT: call void @llvm.memset.p0i8.i64(i8* %[[SS2]], i8 0, i64 24, i32 8, i1 false)
 void test_func_scope_var_local(void) {
   local char *sp1 = 0;
   local char *sp2 = NULL;
   local char *sp3 = (local char*)((void)0, 0);
   const int x = 0;
   local char *sp4 = (local char*)x;
   StructTy1 SS1 = {0, 0, 0, 0, 0};
   StructTy2 SS2 = {0, 0, 0};
 }
 
 
 // Test default initialization of pointers.
 
 // Tentative definition of global variables with non-zero initializer
 // cannot have common linkage since common linkage requires zero initialization
 // and does not have explicit section.
 
 // CHECK: @p1 = common local_unnamed_addr addrspace(1) global i8* null, align 4
 private char *p1;
 
 // CHECK: @p2 = weak local_unnamed_addr addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4
 local char *p2;
 
-// CHECK: @p3 = common local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 4
+// CHECK: @p3 = common local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 8
 constant char *p3;
 
-// CHECK: @p4 = common local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 4
+// CHECK: @p4 = common local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 8
 global char *p4;
 
-// CHECK: @p5 = common local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 4
+// CHECK: @p5 = common local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 8
 generic char *p5;
 
 // Test default initialization of sturcture.
 
-// CHECK: @S1 = weak local_unnamed_addr addrspace(1) global %struct.StructTy1 { i8* null, i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }, align 4
+// CHECK: @S1 = weak local_unnamed_addr addrspace(1) global %struct.StructTy1 { i8* null, i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }, align 8
 StructTy1 S1;
 
-// CHECK: @S2 = common local_unnamed_addr addrspace(1) global %struct.StructTy2 zeroinitializer, align 4
+// CHECK: @S2 = common local_unnamed_addr addrspace(1) global %struct.StructTy2 zeroinitializer, align 8
 StructTy2 S2;
 
 // Test default initialization of array.
-// CHECK: @A1 = weak local_unnamed_addr addrspace(1) global [2 x %struct.StructTy1] [%struct.StructTy1 { i8* null, i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }, %struct.StructTy1 { i8* null, i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }], align 4
+// CHECK: @A1 = weak local_unnamed_addr addrspace(1) global [2 x %struct.StructTy1] [%struct.StructTy1 { i8* null, i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }, %struct.StructTy1 { i8* null, i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8 addrspace(4)* null }], align 8
 StructTy1 A1[2];
 
-// CHECK: @A2 = common local_unnamed_addr addrspace(1) global [2 x %struct.StructTy2] zeroinitializer, align 4
+// CHECK: @A2 = common local_unnamed_addr addrspace(1) global [2 x %struct.StructTy2] zeroinitializer, align 8
 StructTy2 A2[2];
 
 // Test comparison with 0.
 
 // CHECK-LABEL: cmp_private
 // CHECK: icmp eq i8* %p, null
 void cmp_private(private char* p) {
   if (p != 0)

 // CHECK: ret i32 %[[land_ext]]
 int test_and_ptr(private char* p1, local char* p2) {
   return p1 && p2;
 }
 
 // Test folding of null pointer in function scope.
 // NOOPT-LABEL: test_fold_private
 // NOOPT: call void @test_fold_callee
-// NOOPT: store i32 addrspace(1)* null, i32 addrspace(1)** %glob, align 4
+// NOOPT: store i32 addrspace(1)* null, i32 addrspace(1)** %glob, align 8
 // NOOPT: %{{.*}} = sub i64 %{{.*}}, 0
 // NOOPT: call void @test_fold_callee
 // NOOPT: %{{.*}} = add nsw i64 %{{.*}}, 0
 // NOOPT: %{{.*}} = sub nsw i64 %{{.*}}, 1
 void test_fold_callee(void);
 void test_fold_private(void) {
   global int* glob = (test_fold_callee(), (global int*)(generic char*)0);
   long x = glob - (global int*)(generic char*)0;
   x = x + (int)(test_fold_callee(), (private int*)(generic char*)(global short*)0);
   x = x - (int)((private int*)0 == (private int*)(generic char*)0);
 }
 
 // NOOPT-LABEL: test_fold_local
 // NOOPT: call void @test_fold_callee
-// NOOPT: store i32 addrspace(1)* null, i32 addrspace(1)** %glob, align 4
+// NOOPT: store i32 addrspace(1)* null, i32 addrspace(1)** %glob, align 8
 // NOOPT: %{{.*}} = sub i64 %{{.*}}, 0
 // NOOPT: call void @test_fold_callee
 // NOOPT: %{{.*}} = add nsw i64 %{{.*}}, sext (i32 ptrtoint (i32 addrspace(3)* addrspacecast (i32 addrspace(4)* null to i32 addrspace(3)*) to i32) to i64)
 // NOOPT: %{{.*}} = sub nsw i64 %{{.*}}, 1
 void test_fold_local(void) {
   global int* glob = (test_fold_callee(), (global int*)(generic char*)0);
   long x = glob - (global int*)(generic char*)0;
   x = x + (int)(test_fold_callee(), (local int*)(generic char*)(global short*)0);
   x = x - (int)((local int*)0 == (local int*)(generic char*)0);
 }

test/CodeGenOpenCL/amdgpu-sizeof-alignof.cl

+// RUN: %clang_cc1 -triple r600 -cl-std=CL1.2 %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-mesa-mesa3d -cl-std=CL1.2 %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn---opencl -cl-std=CL1.2 %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn---opencl -cl-std=CL2.0 %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn---amdgizcl -cl-std=CL1.2 %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn---amdgizcl -cl-std=CL2.0 %s -emit-llvm -o - | FileCheck %s
+
+#ifdef __AMDGCN__
+#define PTSIZE 8
+#else
+#define PTSIZE 4
+#endif
+
+#ifdef cl_khr_fp64
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+#endif
+#ifdef cl_khr_fp16
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#endif
+
+typedef __SIZE_TYPE__ size_t;
+typedef __PTRDIFF_TYPE__ ptrdiff_t;
+typedef __INTPTR_TYPE__ intptr_t;
+typedef __UINTPTR_TYPE__ uintptr_t;
+typedef global void *global_ptr_t;
+typedef constant void *constant_ptr_t;
+typedef local void *local_ptr_t;
+typedef private void *private_ptr_t;
+
+void check(bool);
+
+void test() {
+  // CHECK-NOT: call void @check(i1 zeroext false)
+  check(sizeof(size_t) == PTSIZE);
+  check(__alignof__(size_t) == PTSIZE);
+  check(sizeof(intptr_t) == PTSIZE);
+  check(__alignof__(intptr_t) == PTSIZE);
+  check(sizeof(uintptr_t) == PTSIZE);
+  check(__alignof__(uintptr_t) == PTSIZE);
+  check(sizeof(ptrdiff_t) == PTSIZE);
+  check(__alignof__(ptrdiff_t) == PTSIZE);
+
+  check(sizeof(char) == 1);
+  check(__alignof__(char) == 1);
+  check(sizeof(short) == 2);
+  check(__alignof__(short) == 2);
+  check(sizeof(int) == 4);
+  check(__alignof__(int) == 4);
+  check(sizeof(long) == 8);
+  check(__alignof__(long) == 8);
+#ifdef cl_khr_fp16
+  check(sizeof(half) == 2);
+  check(__alignof__(half) == 2);
+#endif
+  check(sizeof(float) == 4);
+  check(__alignof__(float) == 4);
+#ifdef cl_khr_fp64
+  check(sizeof(double) == 8);
+  check(__alignof__(double) == 8);
+#endif
+
+  check(sizeof(void*) == (__OPENCL_C_VERSION__ >= 200 ? 8 : 4));
+  check(__alignof__(void*) == (__OPENCL_C_VERSION__ >= 200 ? 8 : 4));
+  check(sizeof(global_ptr_t) == PTSIZE);
+  check(__alignof__(global_ptr_t) == PTSIZE);
+  check(sizeof(constant_ptr_t) == PTSIZE);
+  check(__alignof__(constant_ptr_t) == PTSIZE);
+  check(sizeof(local_ptr_t) == 4);
+  check(__alignof__(private_ptr_t) == 4);
+}

test/Index/pipe-size.cl

 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple x86_64-unknown-linux-gnu %s -o - | FileCheck %s --check-prefix=X86
 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR
 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR64
 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple amdgcn-amd-amdhsa-amdgizcl %s -o - | FileCheck %s --check-prefix=AMD
 __kernel void testPipe( pipe int test )
 {
     int s = sizeof(test);
     // X86: store %opencl.pipe_t* %test, %opencl.pipe_t** %test.addr, align 8
     // X86: store i32 8, i32* %s, align 4
     // SPIR: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)** %test.addr, align 4
     // SPIR: store i32 4, i32* %s, align 4
     // SPIR64: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)** %test.addr, align 8
     // SPIR64: store i32 8, i32* %s, align 4
-    // AMD: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)* addrspace(5)* %test.addr, align 4
+    // AMD: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)* addrspace(5)* %test.addr, align 8
     // AMD: store i32 8, i32 addrspace(5)* %s, align 4
 }

test/SemaCXX/amdgpu-sizeof-alignof.cpp

+// RUN: %clang_cc1 -triple amdgcn---amdgiz -std=c++11 -fsyntax-only -verify %s
+// expected-no-diagnostics
+typedef __SIZE_TYPE__ size_t;
+typedef __PTRDIFF_TYPE__ ptrdiff_t;
+typedef __INTPTR_TYPE__ intptr_t;
+typedef __UINTPTR_TYPE__ uintptr_t;
+typedef __attribute__((address_space(1))) void *global_ptr_t;
+typedef __attribute__((address_space(2))) void *constant_ptr_t;
+typedef __attribute__((address_space(3))) void *local_ptr_t;
+typedef __attribute__((address_space(5))) void *private_ptr_t;
+
+void test() {
+  static_assert(sizeof(size_t) == 8, "bad size");
+  static_assert(alignof(size_t) == 8, "bad alignment");
+  static_assert(sizeof(intptr_t) == 8, "bad size");
+  static_assert(alignof(intptr_t) == 8, "bad alignment");
+  static_assert(sizeof(uintptr_t) == 8, "bad size");
+  static_assert(alignof(uintptr_t) == 8, "bad alignment");
+  static_assert(sizeof(ptrdiff_t) == 8, "bad size");
+  static_assert(alignof(ptrdiff_t) == 8, "bad alignment");
+
+  static_assert(sizeof(char) == 1, "bad size");
+  static_assert(alignof(char) == 1, "bad alignment");
+  static_assert(sizeof(short) == 2, "bad size");
+  static_assert(alignof(short) == 2, "bad alignment");
+  static_assert(sizeof(int) == 4, "bad size");
+  static_assert(alignof(int) == 4, "bad alignment");
+  static_assert(sizeof(long) == 8, "bad size");
+  static_assert(alignof(long) == 8, "bad alignment");
+  static_assert(sizeof(long long) == 8, "bad size");
+  static_assert(alignof(long long) == 8, "bad alignment");
+  static_assert(sizeof(float) == 4, "bad size");
+  static_assert(alignof(float) == 4, "bad alignment");
+  static_assert(sizeof(double) == 8, "bad size");
+  static_assert(alignof(double) == 8, "bad alignment");
+
+  static_assert(sizeof(void*) == 8, "bad size");
+  static_assert(alignof(void*) == 8, "bad alignment");
+  static_assert(sizeof(global_ptr_t) == 8, "bad size");
+  static_assert(alignof(global_ptr_t) == 8, "bad alignment");
+  static_assert(sizeof(constant_ptr_t) == 8, "bad size");
+  static_assert(alignof(constant_ptr_t) == 8, "bad alignment");
+  static_assert(sizeof(local_ptr_t) == 4, "bad size");
+  static_assert(alignof(local_ptr_t) == 4, "bad alignment");
+  static_assert(sizeof(private_ptr_t) == 4, "bad size");
+  static_assert(alignof(private_ptr_t) == 4, "bad alignment");
+}