1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
|
// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fdump-record-layouts \
// RUN: -emit-llvm -o %t -xhip %s 2>&1 | FileCheck %s --check-prefix=AST
// RUN: cat %t | FileCheck --check-prefixes=CHECK,HOST %s
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx1100 \
// RUN: -emit-llvm -fdump-record-layouts -aux-triple x86_64-pc-windows-msvc \
// RUN: -o %t -xhip %s | FileCheck %s --check-prefix=AST
// RUN: cat %t | FileCheck --check-prefixes=CHECK,DEV %s
#include "Inputs/cuda.h"
// AST: *** Dumping AST Record Layout
// AST-LABEL: 0 | struct C
// AST-NEXT: 0 | struct A (base) (empty)
// AST-NEXT: 1 | struct B (base) (empty)
// AST-NEXT: 4 | int i
// AST-NEXT: | [sizeof=8, align=4,
// AST-NEXT: | nvsize=8, nvalign=4]
// CHECK: %struct.C = type { [4 x i8], i32 }
struct A {};
struct B {};
struct C : A, B {
int i;
};
// AST: *** Dumping AST Record Layout
// AST-LABEL: 0 | struct I
// AST-NEXT: 0 | (I vftable pointer)
// AST-NEXT: 8 | int i
// AST-NEXT: | [sizeof=16, align=8,
// AST-NEXT: | nvsize=16, nvalign=8]
// AST: *** Dumping AST Record Layout
// AST-LABEL: 0 | struct J
// AST-NEXT: 0 | struct I (primary base)
// AST-NEXT: 0 | (I vftable pointer)
// AST-NEXT: 8 | int i
// AST-NEXT: 16 | int j
// AST-NEXT: | [sizeof=24, align=8,
// AST-NEXT: | nvsize=24, nvalign=8]
// CHECK: %struct.I = type { ptr, i32 }
// CHECK: %struct.J = type { %struct.I, i32 }
// HOST: @0 = private unnamed_addr constant { [4 x ptr] } { [4 x ptr] [ptr @"??_R4J@@6B@", ptr @"?f@J@@UEAAXXZ", ptr null, ptr @"?h@J@@UEAAXXZ"] }, comdat($"??_7J@@6B@")
// HOST: @1 = private unnamed_addr constant { [4 x ptr] } { [4 x ptr] [ptr @"??_R4I@@6B@", ptr @_purecall, ptr null, ptr @_purecall] }, comdat($"??_7I@@6B@")
// HOST: @"??_7J@@6B@" = unnamed_addr alias ptr, getelementptr inbounds ({ [4 x ptr] }, ptr @0, i32 0, i32 0, i32 1)
// HOST: @"??_7I@@6B@" = unnamed_addr alias ptr, getelementptr inbounds ({ [4 x ptr] }, ptr @1, i32 0, i32 0, i32 1)
// DEV: @_ZTV1J = linkonce_odr unnamed_addr addrspace(1) constant { [5 x ptr addrspace(1)] } { [5 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr @_ZN1J1gEv to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @_ZN1J1hEv to ptr addrspace(1))] }, comdat, align 8
// DEV: @_ZTV1I = linkonce_odr unnamed_addr addrspace(1) constant { [5 x ptr addrspace(1)] } { [5 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr @__cxa_pure_virtual to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @__cxa_pure_virtual to ptr addrspace(1))] }, comdat, align 8
struct I {
virtual void f() = 0;
__device__ virtual void g() = 0;
__device__ __host__ virtual void h() = 0;
int i;
};
struct J : I {
void f() override {}
__device__ void g() override {}
__device__ __host__ void h() override {}
int j;
};
// DEV: define dso_local amdgpu_kernel void @_Z8C_kernel1C(ptr addrspace(4) noundef byref(%struct.C) align 4 %0)
// DEV: %coerce = alloca %struct.C, align 4, addrspace(5)
// DEV: %c = addrspacecast ptr addrspace(5) %coerce to ptr
// DEV: call void @llvm.memcpy.p0.p4.i64(ptr align 4 %c, ptr addrspace(4) align 4 %0, i64 8, i1 false)
// DEV: %i = getelementptr inbounds nuw %struct.C, ptr %c, i32 0, i32 1
// DEV: store i32 1, ptr %i, align 4
__global__ void C_kernel(C c)
{
c.i = 1;
}
// HOST-LABEL: define dso_local void @"?test_C@@YAXXZ"()
// HOST: %c = alloca %struct.C, align 4
// HOST: %i = getelementptr inbounds nuw %struct.C, ptr %c, i32 0, i32 1
// HOST: store i32 11, ptr %i, align 4
void test_C() {
C c;
c.i = 11;
C_kernel<<<1, 1>>>(c);
}
// DEV: define dso_local void @_Z5J_devP1J(ptr noundef %j)
// DEV: %j.addr = alloca ptr, align 8, addrspace(5)
// DEV: %j.addr.ascast = addrspacecast ptr addrspace(5) %j.addr to ptr
// DEV: store ptr %j, ptr %j.addr.ascast, align 8
// DEV: %0 = load ptr, ptr %j.addr.ascast, align 8
// DEV: %i = getelementptr inbounds nuw %struct.I, ptr %0, i32 0, i32 1
// DEV: store i32 2, ptr %i, align 8
// DEV: %1 = load ptr, ptr %j.addr.ascast, align 8
// DEV: %j1 = getelementptr inbounds nuw %struct.J, ptr %1, i32 0, i32 1
// DEV: store i32 3, ptr %j1, align 8
// DEV: %2 = load ptr, ptr %j.addr.ascast, align 8
// DEV: %vtable = load ptr addrspace(1), ptr %2, align 8
// DEV: %vfn = getelementptr inbounds ptr addrspace(1), ptr addrspace(1) %vtable, i64 1
// DEV: %3 = load ptr addrspace(1), ptr addrspace(1) %vfn, align 8
// DEV: call addrspace(1) void %3(ptr noundef nonnull align 8 dereferenceable(24) %2)
// DEV: %4 = load ptr, ptr %j.addr.ascast, align 8
// DEV: %vtable2 = load ptr addrspace(1), ptr %4, align 8
// DEV: %vfn3 = getelementptr inbounds ptr addrspace(1), ptr addrspace(1) %vtable2, i64 2
// DEV: %5 = load ptr addrspace(1), ptr addrspace(1) %vfn3, align 8
// DEV: call addrspace(1) void %5(ptr noundef nonnull align 8 dereferenceable(24) %4)
__device__ void J_dev(J *j) {
j->i = 2;
j->j = 3;
j->g();
j->h();
}
// DEV: define dso_local amdgpu_kernel void @_Z8J_kernelv()
// DEV: %j = alloca %struct.J, align 8, addrspace(5)
// DEV: %j.ascast = addrspacecast ptr addrspace(5) %j to ptr
// DEV: call void @_ZN1JC1Ev(ptr noundef nonnull align 8 dereferenceable(24) %j.ascast)
// DEV: call void @_Z5J_devP1J(ptr noundef %j.ascast)
__global__ void J_kernel() {
J j;
J_dev(&j);
}
// HOST-LABEL: define dso_local void @"?J_host@@YAXPEAUJ@@@Z"(ptr noundef %j)
// HOST: %0 = load ptr, ptr %j.addr, align 8
// HOST: %i = getelementptr inbounds nuw %struct.I, ptr %0, i32 0, i32 1
// HOST: store i32 12, ptr %i, align 8
// HOST: %1 = load ptr, ptr %j.addr, align 8
// HOST: %j1 = getelementptr inbounds nuw %struct.J, ptr %1, i32 0, i32 1
// HOST: store i32 13, ptr %j1, align 8
// HOST: %2 = load ptr, ptr %j.addr, align 8
// HOST: %vtable = load ptr, ptr %2, align 8
// HOST: %vfn = getelementptr inbounds ptr, ptr %vtable, i64 0
// HOST: %3 = load ptr, ptr %vfn, align 8
// HOST: call void %3(ptr noundef nonnull align 8 dereferenceable(24) %2)
// HOST: %4 = load ptr, ptr %j.addr, align 8
// HOST: %vtable2 = load ptr, ptr %4, align 8
// HOST: %vfn3 = getelementptr inbounds ptr, ptr %vtable2, i64 2
// HOST: %5 = load ptr, ptr %vfn3, align 8
// HOST: call void %5(ptr noundef nonnull align 8 dereferenceable(24) %4)
void J_host(J *j) {
j->i = 12;
j->j = 13;
j->f();
j->h();
}
// HOST: define dso_local void @"?test_J@@YAXXZ"()
// HOST: %j = alloca %struct.J, align 8
// HOST: %call = call noundef ptr @"??0J@@QEAA@XZ"(ptr noundef nonnull align 8 dereferenceable(24) %j)
// HOST: call void @"?J_host@@YAXPEAUJ@@@Z"(ptr noundef %j)
void test_J() {
J j;
J_host(&j);
J_kernel<<<1, 1>>>();
}
// HOST: define linkonce_odr dso_local noundef ptr @"??0J@@QEAA@XZ"(ptr noundef nonnull returned align 8 dereferenceable(24) %this)
// HOST: %this.addr = alloca ptr, align 8
// HOST: store ptr %this, ptr %this.addr, align 8
// HOST: %this1 = load ptr, ptr %this.addr, align 8
// HOST: %call = call noundef ptr @"??0I@@QEAA@XZ"(ptr noundef nonnull align 8 dereferenceable(16) %this1) #5
// HOST: store ptr @"??_7J@@6B@", ptr %this1, align 8
// HOST: ret ptr %this1
// HOST: define linkonce_odr dso_local noundef ptr @"??0I@@QEAA@XZ"(ptr noundef nonnull returned align 8 dereferenceable(16) %this)
// HOST: %this.addr = alloca ptr, align 8
// HOST: store ptr %this, ptr %this.addr, align 8
// HOST: %this1 = load ptr, ptr %this.addr, align 8
// HOST: store ptr @"??_7I@@6B@", ptr %this1, align 8
// HOST: ret ptr %this1
// DEV: define linkonce_odr void @_ZN1JC1Ev(ptr noundef nonnull align 8 dereferenceable(24) %this)
// DEV: %this.addr = alloca ptr, align 8, addrspace(5)
// DEV: %this.addr.ascast = addrspacecast ptr addrspace(5) %this.addr to ptr
// DEV: store ptr %this, ptr %this.addr.ascast, align 8
// DEV: %this1 = load ptr, ptr %this.addr.ascast, align 8
// DEV: call void @_ZN1JC2Ev(ptr noundef nonnull align 8 dereferenceable(24) %this1)
// DEV: define linkonce_odr void @_ZN1JC2Ev(ptr noundef nonnull align 8 dereferenceable(24) %this)
// DEV: %this.addr = alloca ptr, align 8, addrspace(5)
// DEV: %this.addr.ascast = addrspacecast ptr addrspace(5) %this.addr to ptr
// DEV: store ptr %this, ptr %this.addr.ascast, align 8
// DEV: %this1 = load ptr, ptr %this.addr.ascast, align 8
// DEV: call void @_ZN1IC2Ev(ptr noundef nonnull align 8 dereferenceable(16) %this1)
// DEV: store ptr addrspace(1) getelementptr inbounds inrange(-16, 24) ({ [5 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTV1J, i32 0, i32 0, i32 2), ptr %this1, align 8
// DEV: define linkonce_odr void @_ZN1IC2Ev(ptr noundef nonnull align 8 dereferenceable(16) %this)
// DEV: %this.addr = alloca ptr, align 8, addrspace(5)
// DEV: %this.addr.ascast = addrspacecast ptr addrspace(5) %this.addr to ptr
// DEV: store ptr %this, ptr %this.addr.ascast, align 8
// DEV: %this1 = load ptr, ptr %this.addr.ascast, align 8
// DEV: store ptr addrspace(1) getelementptr inbounds inrange(-16, 24) ({ [5 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTV1I, i32 0, i32 0, i32 2), ptr %this1, align 8
|
