Skip to content

Commit 25d1b43

Browse files
yxsamliuyxsamliu
committedJul 5, 2017
[AMDGPU] Fix size and alignment of size_t and pointer types
Differential Revision: https://reviews.llvm.org/D34995 llvm-svn: 307121
1 parent 845a796 commit 25d1b43

8 files changed

+167
-40
lines changed
 

‎clang/lib/Basic/Targets.cpp

+14-6
Original file line numberDiff line numberDiff line change
@@ -2189,6 +2189,15 @@ class AMDGPUTargetInfo final : public TargetInfo {
21892189
Triple.getEnvironmentName() == "amdgizcl" ||
21902190
!isAMDGCN(Triple));
21912191
UseAddrSpaceMapMangling = true;
2192+
2193+
// Set pointer width and alignment for target address space 0.
2194+
PointerWidth = PointerAlign = DataLayout->getPointerSizeInBits();
2195+
if (getMaxPointerWidth() == 64) {
2196+
LongWidth = LongAlign = 64;
2197+
SizeType = UnsignedLong;
2198+
PtrDiffType = SignedLong;
2199+
IntPtrType = SignedLong;
2200+
}
21922201
}
21932202

21942203
void setAddressSpaceMap(bool DefaultIsPrivate) {
@@ -2216,6 +2225,10 @@ class AMDGPUTargetInfo final : public TargetInfo {
22162225
return 64;
22172226
}
22182227

2228+
uint64_t getPointerAlignV(unsigned AddrSpace) const override {
2229+
return getPointerWidthV(AddrSpace);
2230+
}
2231+
22192232
uint64_t getMaxPointerWidth() const override {
22202233
return getTriple().getArch() == llvm::Triple::amdgcn ? 64 : 32;
22212234
}
@@ -2392,12 +2405,7 @@ class AMDGPUTargetInfo final : public TargetInfo {
23922405
}
23932406

23942407
/// \returns Target specific vtbl ptr address space.
2395-
unsigned getVtblPtrAddressSpace() const override {
2396-
// \todo: We currently have address spaces defined in AMDGPU Backend. It
2397-
// would be nice if we could use it here instead of using bare numbers (same
2398-
// applies to getDWARFAddressSpace).
2399-
return 2; // constant.
2400-
}
2408+
unsigned getVtblPtrAddressSpace() const override { return AS.Constant; }
24012409

24022410
/// \returns If a target requires an address within a target specific address
24032411
/// space \p AddressSpace to be converted in order to be used, then return the

‎clang/test/CodeGen/default-address-space.c

+1-1
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,7 @@ void test3() {
5050
// PIZ: %[[arrayidx:.*]] = getelementptr inbounds i32, i32 addrspace(4)* %[[r0]]
5151
// PIZ: store i32 0, i32 addrspace(4)* %[[arrayidx]]
5252
// CHECK-LABEL: define void @test4(i32* %a)
53-
// CHECK: %[[alloca:.*]] = alloca i32*, align 4, addrspace(5)
53+
// CHECK: %[[alloca:.*]] = alloca i32*, align 8, addrspace(5)
5454
// CHECK: %[[a_addr:.*]] = addrspacecast{{.*}} %[[alloca]] to i32**
5555
// CHECK: store i32* %a, i32** %[[a_addr]]
5656
// CHECK: %[[r0:.*]] = load i32*, i32** %[[a_addr]]

‎clang/test/CodeGenCXX/amdgcn-automatic-variable.cpp

+6-4
Original file line numberDiff line numberDiff line change
@@ -15,8 +15,8 @@ void func2(void) {
1515
// CHECK: %lv1 = alloca i32, align 4, addrspace(5)
1616
// CHECK: %lv2 = alloca i32, align 4, addrspace(5)
1717
// CHECK: %la = alloca [100 x i32], align 4, addrspace(5)
18-
// CHECK: %lp1 = alloca i32*, align 4, addrspace(5)
19-
// CHECK: %lp2 = alloca i32*, align 4, addrspace(5)
18+
// CHECK: %lp1 = alloca i32*, align 8, addrspace(5)
19+
// CHECK: %lp2 = alloca i32*, align 8, addrspace(5)
2020
// CHECK: %lvc = alloca i32, align 4, addrspace(5)
2121

2222
// CHECK: %[[r0:.*]] = addrspacecast i32 addrspace(5)* %lv1 to i32*
@@ -34,12 +34,12 @@ void func2(void) {
3434
la[0] = 3;
3535

3636
// CHECK: %[[r3:.*]] = addrspacecast i32* addrspace(5)* %lp1 to i32**
37-
// CHECK: store i32* %[[r0]], i32** %[[r3]], align 4
37+
// CHECK: store i32* %[[r0]], i32** %[[r3]], align 8
3838
int *lp1 = &lv1;
3939

4040
// CHECK: %[[r4:.*]] = addrspacecast i32* addrspace(5)* %lp2 to i32**
4141
// CHECK: %[[arraydecay:.*]] = getelementptr inbounds [100 x i32], [100 x i32]* %[[r2]], i32 0, i32 0
42-
// CHECK: store i32* %[[arraydecay]], i32** %[[r4]], align 4
42+
// CHECK: store i32* %[[arraydecay]], i32** %[[r4]], align 8
4343
int *lp2 = la;
4444

4545
// CHECK: call void @_Z5func1Pi(i32* %[[r0]])
@@ -80,3 +80,5 @@ void func4(int x) {
8080
// CHECK: call void @_Z5func1Pi(i32* %[[r0]])
8181
func1(&x);
8282
}
83+
84+
// CHECK-NOT: !opencl.ocl.version

‎clang/test/CodeGenOpenCL/amdgcn-automatic-variable.cl

+4-4
Original file line numberDiff line numberDiff line change
@@ -22,8 +22,8 @@ void func2(void) {
2222
// CHECK: %la = alloca [100 x i32], align 4, addrspace(5)
2323
// CL12: %lp1 = alloca i32 addrspace(5)*, align 4, addrspace(5)
2424
// CL12: %lp2 = alloca i32 addrspace(5)*, align 4, addrspace(5)
25-
// CL20: %lp1 = alloca i32*, align 4, addrspace(5)
26-
// CL20: %lp2 = alloca i32*, align 4, addrspace(5)
25+
// CL20: %lp1 = alloca i32*, align 8, addrspace(5)
26+
// CL20: %lp2 = alloca i32*, align 8, addrspace(5)
2727
// CHECK: %lvc = alloca i32, align 4, addrspace(5)
2828

2929
// CHECK: store i32 1, i32 addrspace(5)* %lv1
@@ -39,13 +39,13 @@ void func2(void) {
3939

4040
// CL12: store i32 addrspace(5)* %lv1, i32 addrspace(5)* addrspace(5)* %lp1, align 4
4141
// CL20: %[[r0:.*]] = addrspacecast i32 addrspace(5)* %lv1 to i32*
42-
// CL20: store i32* %[[r0]], i32* addrspace(5)* %lp1, align 4
42+
// CL20: store i32* %[[r0]], i32* addrspace(5)* %lp1, align 8
4343
int *lp1 = &lv1;
4444

4545
// CHECK: %[[arraydecay:.*]] = getelementptr inbounds [100 x i32], [100 x i32] addrspace(5)* %la, i32 0, i32 0
4646
// CL12: store i32 addrspace(5)* %[[arraydecay]], i32 addrspace(5)* addrspace(5)* %lp2, align 4
4747
// CL20: %[[r1:.*]] = addrspacecast i32 addrspace(5)* %[[arraydecay]] to i32*
48-
// CL20: store i32* %[[r1]], i32* addrspace(5)* %lp2, align 4
48+
// CL20: store i32* %[[r1]], i32* addrspace(5)* %lp2, align 8
4949
int *lp2 = la;
5050

5151
// CL12: call void @func1(i32 addrspace(5)* %lv1)

‎clang/test/CodeGenOpenCL/amdgpu-nullptr.cl

+24-24
Original file line numberDiff line numberDiff line change
@@ -27,13 +27,13 @@ private char *private_p = 0;
2727
// CHECK: @local_p = local_unnamed_addr addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4
2828
local char *local_p = 0;
2929

30-
// CHECK: @global_p = local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 4
30+
// CHECK: @global_p = local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 8
3131
global char *global_p = 0;
3232

33-
// CHECK: @constant_p = local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 4
33+
// CHECK: @constant_p = local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 8
3434
constant char *constant_p = 0;
3535

36-
// CHECK: @generic_p = local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 4
36+
// CHECK: @generic_p = local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 8
3737
generic char *generic_p = 0;
3838

3939
// Test NULL as initializer.
@@ -44,19 +44,19 @@ private char *private_p_NULL = NULL;
4444
// CHECK: @local_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4
4545
local char *local_p_NULL = NULL;
4646

47-
// CHECK: @global_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 4
47+
// CHECK: @global_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 8
4848
global char *global_p_NULL = NULL;
4949

50-
// CHECK: @constant_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 4
50+
// CHECK: @constant_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 8
5151
constant char *constant_p_NULL = NULL;
5252

53-
// CHECK: @generic_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 4
53+
// CHECK: @generic_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 8
5454
generic char *generic_p_NULL = NULL;
5555

5656
// Test constant folding of null pointer.
5757
// A null pointer should be folded to a null pointer in the target address space.
5858

59-
// CHECK: @fold_generic = local_unnamed_addr addrspace(1) global i32 addrspace(4)* null, align 4
59+
// CHECK: @fold_generic = local_unnamed_addr addrspace(1) global i32 addrspace(4)* null, align 8
6060
generic int *fold_generic = (global int*)(generic float*)(private char*)0;
6161

6262
// CHECK: @fold_priv = local_unnamed_addr addrspace(1) global i16* null, align 4
@@ -104,8 +104,8 @@ int fold_int5_local = (int) &((local StructTy1*)0)->p2;
104104
// NOOPT: @test_static_var_private.sp3 = internal addrspace(1) global i8* null, align 4
105105
// NOOPT: @test_static_var_private.sp4 = internal addrspace(1) global i8* null, align 4
106106
// NOOPT: @test_static_var_private.sp5 = internal addrspace(1) global i8* null, align 4
107-
// 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
108-
// NOOPT: @test_static_var_private.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 4
107+
// 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
108+
// NOOPT: @test_static_var_private.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 8
109109

110110
void test_static_var_private(void) {
111111
static private char *sp1 = 0;
@@ -123,8 +123,8 @@ void test_static_var_private(void) {
123123
// NOOPT: @test_static_var_local.sp3 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4
124124
// NOOPT: @test_static_var_local.sp4 = internal addrspace(1) global i8 addrspace(3)* null, align 4
125125
// NOOPT: @test_static_var_local.sp5 = internal addrspace(1) global i8 addrspace(3)* null, align 4
126-
// 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
127-
// NOOPT: @test_static_var_local.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 4
126+
// 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
127+
// NOOPT: @test_static_var_local.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 8
128128
void test_static_var_local(void) {
129129
static local char *sp1 = 0;
130130
static local char *sp2 = NULL;
@@ -143,9 +143,9 @@ void test_static_var_local(void) {
143143
// NOOPT: store i8* null, i8** %sp3, align 4
144144
// NOOPT: store i8* null, i8** %sp4, align 4
145145
// NOOPT: %[[SS1:.*]] = bitcast %struct.StructTy1* %SS1 to i8*
146-
// 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)
146+
// 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)
147147
// NOOPT: %[[SS2:.*]] = bitcast %struct.StructTy2* %SS2 to i8*
148-
// NOOPT: call void @llvm.memset.p0i8.i64(i8* %[[SS2]], i8 0, i64 24, i32 4, i1 false)
148+
// NOOPT: call void @llvm.memset.p0i8.i64(i8* %[[SS2]], i8 0, i64 24, i32 8, i1 false)
149149
void test_func_scope_var_private(void) {
150150
private char *sp1 = 0;
151151
private char *sp2 = NULL;
@@ -163,9 +163,9 @@ void test_func_scope_var_private(void) {
163163
// NOOPT: store i8 addrspace(3)* null, i8 addrspace(3)** %sp3, align 4
164164
// NOOPT: store i8 addrspace(3)* null, i8 addrspace(3)** %sp4, align 4
165165
// NOOPT: %[[SS1:.*]] = bitcast %struct.StructTy1* %SS1 to i8*
166-
// 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)
166+
// 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)
167167
// NOOPT: %[[SS2:.*]] = bitcast %struct.StructTy2* %SS2 to i8*
168-
// NOOPT: call void @llvm.memset.p0i8.i64(i8* %[[SS2]], i8 0, i64 24, i32 4, i1 false)
168+
// NOOPT: call void @llvm.memset.p0i8.i64(i8* %[[SS2]], i8 0, i64 24, i32 8, i1 false)
169169
void test_func_scope_var_local(void) {
170170
local char *sp1 = 0;
171171
local char *sp2 = NULL;
@@ -189,28 +189,28 @@ private char *p1;
189189
// CHECK: @p2 = weak local_unnamed_addr addrspace(1) global i8 addrspace(3)* addrspacecast (i8 addrspace(4)* null to i8 addrspace(3)*), align 4
190190
local char *p2;
191191

192-
// CHECK: @p3 = common local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 4
192+
// CHECK: @p3 = common local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 8
193193
constant char *p3;
194194

195-
// CHECK: @p4 = common local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 4
195+
// CHECK: @p4 = common local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 8
196196
global char *p4;
197197

198-
// CHECK: @p5 = common local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 4
198+
// CHECK: @p5 = common local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 8
199199
generic char *p5;
200200

201201
// Test default initialization of sturcture.
202202

203-
// 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
203+
// 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
204204
StructTy1 S1;
205205

206-
// CHECK: @S2 = common local_unnamed_addr addrspace(1) global %struct.StructTy2 zeroinitializer, align 4
206+
// CHECK: @S2 = common local_unnamed_addr addrspace(1) global %struct.StructTy2 zeroinitializer, align 8
207207
StructTy2 S2;
208208

209209
// Test default initialization of array.
210-
// 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
210+
// 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
211211
StructTy1 A1[2];
212212

213-
// CHECK: @A2 = common local_unnamed_addr addrspace(1) global [2 x %struct.StructTy2] zeroinitializer, align 4
213+
// CHECK: @A2 = common local_unnamed_addr addrspace(1) global [2 x %struct.StructTy2] zeroinitializer, align 8
214214
StructTy2 A2[2];
215215

216216
// Test comparison with 0.
@@ -597,7 +597,7 @@ int test_and_ptr(private char* p1, local char* p2) {
597597
// Test folding of null pointer in function scope.
598598
// NOOPT-LABEL: test_fold_private
599599
// NOOPT: call void @test_fold_callee
600-
// NOOPT: store i32 addrspace(1)* null, i32 addrspace(1)** %glob, align 4
600+
// NOOPT: store i32 addrspace(1)* null, i32 addrspace(1)** %glob, align 8
601601
// NOOPT: %{{.*}} = sub i64 %{{.*}}, 0
602602
// NOOPT: call void @test_fold_callee
603603
// NOOPT: %{{.*}} = add nsw i64 %{{.*}}, 0
@@ -612,7 +612,7 @@ void test_fold_private(void) {
612612

613613
// NOOPT-LABEL: test_fold_local
614614
// NOOPT: call void @test_fold_callee
615-
// NOOPT: store i32 addrspace(1)* null, i32 addrspace(1)** %glob, align 4
615+
// NOOPT: store i32 addrspace(1)* null, i32 addrspace(1)** %glob, align 8
616616
// NOOPT: %{{.*}} = sub i64 %{{.*}}, 0
617617
// NOOPT: call void @test_fold_callee
618618
// NOOPT: %{{.*}} = add nsw i64 %{{.*}}, sext (i32 ptrtoint (i32 addrspace(3)* addrspacecast (i32 addrspace(4)* null to i32 addrspace(3)*) to i32) to i64)

‎clang/test/CodeGenOpenCL/amdgpu-sizeof-alignof.cl

+70
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,70 @@
1+
// RUN: %clang_cc1 -triple r600 -cl-std=CL1.2 %s -emit-llvm -o - | FileCheck %s
2+
// RUN: %clang_cc1 -triple amdgcn-mesa-mesa3d -cl-std=CL1.2 %s -emit-llvm -o - | FileCheck %s
3+
// RUN: %clang_cc1 -triple amdgcn---opencl -cl-std=CL1.2 %s -emit-llvm -o - | FileCheck %s
4+
// RUN: %clang_cc1 -triple amdgcn---opencl -cl-std=CL2.0 %s -emit-llvm -o - | FileCheck %s
5+
// RUN: %clang_cc1 -triple amdgcn---amdgizcl -cl-std=CL1.2 %s -emit-llvm -o - | FileCheck %s
6+
// RUN: %clang_cc1 -triple amdgcn---amdgizcl -cl-std=CL2.0 %s -emit-llvm -o - | FileCheck %s
7+
8+
#ifdef __AMDGCN__
9+
#define PTSIZE 8
10+
#else
11+
#define PTSIZE 4
12+
#endif
13+
14+
#ifdef cl_khr_fp64
15+
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
16+
#endif
17+
#ifdef cl_khr_fp16
18+
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
19+
#endif
20+
21+
typedef __SIZE_TYPE__ size_t;
22+
typedef __PTRDIFF_TYPE__ ptrdiff_t;
23+
typedef __INTPTR_TYPE__ intptr_t;
24+
typedef __UINTPTR_TYPE__ uintptr_t;
25+
typedef global void *global_ptr_t;
26+
typedef constant void *constant_ptr_t;
27+
typedef local void *local_ptr_t;
28+
typedef private void *private_ptr_t;
29+
30+
void check(bool);
31+
32+
void test() {
33+
// CHECK-NOT: call void @check(i1 zeroext false)
34+
check(sizeof(size_t) == PTSIZE);
35+
check(__alignof__(size_t) == PTSIZE);
36+
check(sizeof(intptr_t) == PTSIZE);
37+
check(__alignof__(intptr_t) == PTSIZE);
38+
check(sizeof(uintptr_t) == PTSIZE);
39+
check(__alignof__(uintptr_t) == PTSIZE);
40+
check(sizeof(ptrdiff_t) == PTSIZE);
41+
check(__alignof__(ptrdiff_t) == PTSIZE);
42+
43+
check(sizeof(char) == 1);
44+
check(__alignof__(char) == 1);
45+
check(sizeof(short) == 2);
46+
check(__alignof__(short) == 2);
47+
check(sizeof(int) == 4);
48+
check(__alignof__(int) == 4);
49+
check(sizeof(long) == 8);
50+
check(__alignof__(long) == 8);
51+
#ifdef cl_khr_fp16
52+
check(sizeof(half) == 2);
53+
check(__alignof__(half) == 2);
54+
#endif
55+
check(sizeof(float) == 4);
56+
check(__alignof__(float) == 4);
57+
#ifdef cl_khr_fp64
58+
check(sizeof(double) == 8);
59+
check(__alignof__(double) == 8);
60+
#endif
61+
62+
check(sizeof(void*) == (__OPENCL_C_VERSION__ >= 200 ? 8 : 4));
63+
check(__alignof__(void*) == (__OPENCL_C_VERSION__ >= 200 ? 8 : 4));
64+
check(sizeof(global_ptr_t) == PTSIZE);
65+
check(__alignof__(global_ptr_t) == PTSIZE);
66+
check(sizeof(constant_ptr_t) == PTSIZE);
67+
check(__alignof__(constant_ptr_t) == PTSIZE);
68+
check(sizeof(local_ptr_t) == 4);
69+
check(__alignof__(private_ptr_t) == 4);
70+
}

‎clang/test/Index/pipe-size.cl

+1-1
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,6 @@ __kernel void testPipe( pipe int test )
1111
// SPIR: store i32 4, i32* %s, align 4
1212
// SPIR64: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)** %test.addr, align 8
1313
// SPIR64: store i32 8, i32* %s, align 4
14-
// AMD: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)* addrspace(5)* %test.addr, align 4
14+
// AMD: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)* addrspace(5)* %test.addr, align 8
1515
// AMD: store i32 8, i32 addrspace(5)* %s, align 4
1616
}

‎clang/test/SemaCXX/amdgpu-sizeof-alignof.cpp

+47
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,47 @@
1+
// RUN: %clang_cc1 -triple amdgcn---amdgiz -std=c++11 -fsyntax-only -verify %s
2+
// expected-no-diagnostics
3+
typedef __SIZE_TYPE__ size_t;
4+
typedef __PTRDIFF_TYPE__ ptrdiff_t;
5+
typedef __INTPTR_TYPE__ intptr_t;
6+
typedef __UINTPTR_TYPE__ uintptr_t;
7+
typedef __attribute__((address_space(1))) void *global_ptr_t;
8+
typedef __attribute__((address_space(2))) void *constant_ptr_t;
9+
typedef __attribute__((address_space(3))) void *local_ptr_t;
10+
typedef __attribute__((address_space(5))) void *private_ptr_t;
11+
12+
void test() {
13+
static_assert(sizeof(size_t) == 8, "bad size");
14+
static_assert(alignof(size_t) == 8, "bad alignment");
15+
static_assert(sizeof(intptr_t) == 8, "bad size");
16+
static_assert(alignof(intptr_t) == 8, "bad alignment");
17+
static_assert(sizeof(uintptr_t) == 8, "bad size");
18+
static_assert(alignof(uintptr_t) == 8, "bad alignment");
19+
static_assert(sizeof(ptrdiff_t) == 8, "bad size");
20+
static_assert(alignof(ptrdiff_t) == 8, "bad alignment");
21+
22+
static_assert(sizeof(char) == 1, "bad size");
23+
static_assert(alignof(char) == 1, "bad alignment");
24+
static_assert(sizeof(short) == 2, "bad size");
25+
static_assert(alignof(short) == 2, "bad alignment");
26+
static_assert(sizeof(int) == 4, "bad size");
27+
static_assert(alignof(int) == 4, "bad alignment");
28+
static_assert(sizeof(long) == 8, "bad size");
29+
static_assert(alignof(long) == 8, "bad alignment");
30+
static_assert(sizeof(long long) == 8, "bad size");
31+
static_assert(alignof(long long) == 8, "bad alignment");
32+
static_assert(sizeof(float) == 4, "bad size");
33+
static_assert(alignof(float) == 4, "bad alignment");
34+
static_assert(sizeof(double) == 8, "bad size");
35+
static_assert(alignof(double) == 8, "bad alignment");
36+
37+
static_assert(sizeof(void*) == 8, "bad size");
38+
static_assert(alignof(void*) == 8, "bad alignment");
39+
static_assert(sizeof(global_ptr_t) == 8, "bad size");
40+
static_assert(alignof(global_ptr_t) == 8, "bad alignment");
41+
static_assert(sizeof(constant_ptr_t) == 8, "bad size");
42+
static_assert(alignof(constant_ptr_t) == 8, "bad alignment");
43+
static_assert(sizeof(local_ptr_t) == 4, "bad size");
44+
static_assert(alignof(local_ptr_t) == 4, "bad alignment");
45+
static_assert(sizeof(private_ptr_t) == 4, "bad size");
46+
static_assert(alignof(private_ptr_t) == 4, "bad alignment");
47+
}

0 commit comments

Comments
 (0)
Please sign in to comment.