Skip to content

Commit b00fb7c

Browse files
elizabethandrewselizabethandrews
authored
[SYCL] Fix crash when kernel object field is an array of structs (#2083)
This patch adds a 'top-level' struct InitListExpr for each element of a struct array. This is required to nest the InitListExprs of data members of each struct element correctly under the corresponding struct element. Signed-off-by: Elizabeth Andrews <[email protected]>
1 parent 4dbd081 commit b00fb7c

File tree

3 files changed

+225
-7
lines changed

3 files changed

+225
-7
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1598,9 +1598,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
15981598
}
15991599

16001600
void addStructInit(const CXXRecordDecl *RD) {
1601-
if (!RD)
1602-
return;
1603-
16041601
const ASTRecordLayout &Info =
16051602
SemaRef.getASTContext().getASTRecordLayout(RD);
16061603
int NumberOfFields = Info.getFieldCount();
@@ -1621,7 +1618,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
16211618
}
16221619

16231620
bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD) final {
1624-
const CXXRecordDecl *RD = FD->getType()->getAsCXXRecordDecl();
1621+
// Handle struct when kernel object field is struct type or array of
1622+
// structs.
1623+
const CXXRecordDecl *RD =
1624+
FD->getType()->getBaseElementTypeUnsafe()->getAsCXXRecordDecl();
16251625

16261626
// Initializers for accessors inside stream not added.
16271627
if (!Util::isSyclStreamType(FD->getType()))

‎clang/test/CodeGenSYCL/kernel-param-pod-array.cpp

100755100644
Lines changed: 111 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,14 +11,31 @@ __attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
1111
kernelFunc();
1212
}
1313

14+
struct foo_inner {
15+
int foo_inner_x;
16+
int foo_inner_y;
17+
};
18+
19+
struct foo {
20+
int foo_a;
21+
foo_inner foo_b[2];
22+
int foo_c;
23+
};
24+
1425
int main() {
1526

1627
int a[2];
28+
foo struct_array[2];
1729

1830
a_kernel<class kernel_B>(
1931
[=]() {
2032
int local = a[1];
2133
});
34+
35+
a_kernel<class kernel_C>(
36+
[=]() {
37+
foo local = struct_array[1];
38+
});
2239
}
2340

2441
// Check kernel_B parameters
@@ -40,4 +57,97 @@ int main() {
4057
// CHECK: store i32 [[ARRAY0]], i32* [[ARRAY_BEGIN]], align 4
4158
// CHECK: [[ARRAY_ELEMENT:%[a-zA-Z_.]+]] = getelementptr inbounds i32, i32* %arrayinit.begin, i64 1
4259
// CHECK: [[ARRAY1:%[0-9]*]] = load i32, i32* [[ELEM_L1]], align 4
43-
// CHECK: store i32 [[ARRAY1]], i32* [[ARRAY_ELEMENT]], align 4
60+
// CHECK: store i32 [[ARRAY1]], i32* [[ARRAY_ELEMENT]], align 4
61+
62+
// Check kernel_C parameters
63+
// CHECK: define spir_kernel void @{{.*}}kernel_C
64+
// CHECK-SAME: i32 [[FOO1_A:%[a-zA-Z0-9_]+]], i32 [[FOO1_B1_X:%[a-zA-Z0-9_]+]], i32 [[FOO1_B1_Y:%[a-zA-Z0-9_]+]], i32 [[FOO1_B2_X:%[a-zA-Z0-9_]+]], i32 [[FOO1_B2_Y:%[a-zA-Z0-9_]+]], i32 [[FOO1_C:%[a-zA-Z0-9_]+]],
65+
// CHECK-SAME: i32 [[FOO2_A:%[a-zA-Z0-9_]+]], i32 [[FOO2_B1_X:%[a-zA-Z0-9_]+]], i32 [[FOO2_B1_Y:%[a-zA-Z0-9_]+]], i32 [[FOO2_B2_X:%[a-zA-Z0-9_]+]], i32 [[FOO2_B2_Y:%[a-zA-Z0-9_]+]], i32 [[FOO2_C:%[a-zA-Z0-9_]+]]
66+
67+
// Check local lambda object alloca
68+
// CHECK: [[KERNEL_OBJ:%[0-9]+]] = alloca %"class.{{.*}}.anon.0", align 4
69+
70+
// Check local stores
71+
// CHECK: store i32 [[FOO1_A]], i32* [[FOO1_A_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
72+
// CHECK: store i32 [[FOO1_B1_X]], i32* [[FOO1_B1_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
73+
// CHECK: store i32 [[FOO1_B1_Y]], i32* [[FOO1_B1_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
74+
// CHECK: store i32 [[FOO1_B2_X]], i32* [[FOO1_B2_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
75+
// CHECK: store i32 [[FOO1_B2_Y]], i32* [[FOO1_B2_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
76+
// CHECK: store i32 [[FOO1_C]], i32* [[FOO1_C_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
77+
// CHECK: store i32 [[FOO2_A]], i32* [[FOO2_A_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
78+
// CHECK: store i32 [[FOO2_B1_X]], i32* [[FOO2_B1_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
79+
// CHECK: store i32 [[FOO2_B1_Y]], i32* [[FOO2_B1_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
80+
// CHECK: store i32 [[FOO2_B2_X]], i32* [[FOO2_B2_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
81+
// CHECK: store i32 [[FOO2_B2_Y]], i32* [[FOO2_B2_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
82+
// CHECK: store i32 [[FOO2_C]], i32* [[FOO2_C_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
83+
84+
// Check initialization of local array
85+
86+
// Initialize struct_array[0].foo_a
87+
// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon.0", %"class.{{.*}}.anon.0"* [[KERNEL_OBJ]], i32 0, i32 0
88+
// CHECK: [[FOO_ARRAY_0:%[a-zA-Z_.]+]] = getelementptr inbounds [2 x %struct.{{.*}}.foo], [2 x %struct.{{.*}}.foo]* [[GEP]], i64 0, i64 0
89+
// CHECK: [[GEP_FOO1_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* [[FOO_ARRAY_0]], i32 0, i32 0
90+
// CHECK: [[LOAD_FOO1_A:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_A_LOCAL]], align 4
91+
// CHECK: store i32 [[LOAD_FOO1_A]], i32* [[GEP_FOO1_A]], align 4
92+
93+
// Initialize struct_array[0].foo_b[0].x
94+
// CHECK: [[GEP_FOO1_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* [[FOO_ARRAY_0]], i32 0, i32 1
95+
// CHECK: [[B_ARRAY_0:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [2 x %struct.{{.*}}foo_inner.foo_inner], [2 x %struct.{{.*}}foo_inner.foo_inner]* [[GEP_FOO1_B]], i64 0, i64 0
96+
// CHECK: [[GEP_FOO1_B1_X:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_0]], i32 0, i32 0
97+
// CHECK: [[LOAD_FOO1_B1_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B1_X_LOCAL]], align 4
98+
// CHECK: store i32 [[LOAD_FOO1_B1_X]], i32* [[GEP_FOO1_B1_X]], align 4
99+
100+
// Initialize struct_array[0].foo_b[0].y
101+
// CHECK: [[GEP_FOO1_B1_Y:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_0]], i32 0, i32 1
102+
// CHECK: [[LOAD_FOO1_B1_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B1_Y_LOCAL]], align 4
103+
// CHECK: store i32 [[LOAD_FOO1_B1_Y]], i32* [[GEP_FOO1_B1_Y]], align 4
104+
105+
// Initialize struct_array[0].foo_b[1].x
106+
// CHECK: [[B_ARRAY_1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_0]], i64 1
107+
// CHECK: [[GEP_FOO1_B2_X:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_1]], i32 0, i32 0
108+
// CHECK: [[LOAD_FOO1_B2_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B2_X_LOCAL]], align 4
109+
// CHECK: store i32 [[LOAD_FOO1_B2_X]], i32* [[GEP_FOO1_B2_X]], align 4
110+
111+
// Initialize struct_array[0].foo_b[1].y
112+
// CHECK: [[GEP_FOO1_B2_Y:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_1]], i32 0, i32 1
113+
// CHECK: [[LOAD_FOO1_B2_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B2_Y_LOCAL]], align 4
114+
// CHECK: store i32 [[LOAD_FOO1_B2_Y]], i32* [[GEP_FOO1_B2_Y]], align 4
115+
116+
// Initialize struct_array[0].foo_c
117+
// CHECK: [[GEP_FOO1_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_0]], i32 0, i32 2
118+
// CHECK: [[LOAD_FOO1_C:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_C_LOCAL]], align 4
119+
// CHECK: store i32 [[LOAD_FOO1_C]], i32* [[GEP_FOO1_C]], align 4
120+
121+
// Initialize struct_array[1].foo_a
122+
// CHECK: [[FOO_ARRAY_1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct._ZTS3foo.foo, %struct._ZTS3foo.foo* [[FOO_ARRAY_0]], i64 1
123+
// CHECK: [[GEP_FOO2_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_1]], i32 0, i32 0
124+
// CHECK: [[LOAD_FOO2_A:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_A_LOCAL]], align 4
125+
// CHECK: store i32 [[LOAD_FOO2_A]], i32* [[GEP_FOO2_A]], align 4
126+
127+
// Initialize struct_array[1].foo_b[0].x
128+
// CHECK: [[GEP_FOO2_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* [[FOO_ARRAY_1]], i32 0, i32 1
129+
// CHECK: [[FOO2_B_ARRAY_0:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [2 x %struct.{{.*}}foo_inner.foo_inner], [2 x %struct.{{.*}}foo_inner.foo_inner]* [[GEP_FOO2_B]], i64 0, i64 0
130+
// CHECK: [[GEP_FOO2_B1_X:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_0]], i32 0, i32 0
131+
// CHECK: [[LOAD_FOO2_B1_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B1_X_LOCAL]], align 4
132+
// CHECK: store i32 [[LOAD_FOO2_B1_X]], i32* [[GEP_FOO2_B1_X]]
133+
134+
// Initialize struct_array[1].foo_b[0].y
135+
// CHECK: [[GEP_FOO2_B1_Y:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_0]], i32 0, i32 1
136+
// CHECK: [[LOAD_FOO2_B1_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B1_Y_LOCAL]], align 4
137+
// CHECK: store i32 [[LOAD_FOO2_B1_Y]], i32* [[GEP_FOO2_B1_Y]], align 4
138+
139+
// Initialize struct_array[1].foo_b[1].x
140+
// CHECK: [[FOO2_B_ARRAY_1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_0]], i64 1
141+
// CHECK: [[GEP_FOO2_B2_X:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_1]], i32 0, i32 0
142+
// CHECK: [[LOAD_FOO2_B2_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B2_X_LOCAL]], align 4
143+
// CHECK: store i32 [[LOAD_FOO2_B2_X]], i32* [[GEP_FOO2_B2_X]], align 4
144+
145+
// Initialize struct_array[1].foo_b[1].y
146+
// CHECK: [[GEP_FOO2_B2_Y:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_1]], i32 0, i32 1
147+
// CHECK: [[LOAD_FOO2_B2_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B2_Y_LOCAL]], align 4
148+
// CHECK: store i32 [[LOAD_FOO2_B2_Y]], i32* [[GEP_FOO2_B2_Y]], align 4
149+
150+
// Initialize struct_array[1].foo_c
151+
// CHECK: [[GEP_FOO2_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_1]], i32 0, i32 2
152+
// CHECK: [[LOAD_FOO2_C:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_C_LOCAL]], align 4
153+
// CHECK: store i32 [[LOAD_FOO2_C]], i32* [[GEP_FOO2_C]], align 4

‎clang/test/SemaSYCL/array-kernel-param.cpp

Lines changed: 110 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -23,6 +23,20 @@ int main() {
2323
Accessor member_acc[2];
2424
} struct_acc;
2525

26+
struct foo_inner {
27+
int foo_inner_x;
28+
int foo_inner_y;
29+
int foo_inner_z[2];
30+
};
31+
32+
struct foo {
33+
int foo_a;
34+
foo_inner foo_b[2];
35+
int foo_c;
36+
};
37+
38+
foo struct_array[2];
39+
2640
a_kernel<class kernel_A>(
2741
[=]() {
2842
acc[1].use();
@@ -37,6 +51,11 @@ int main() {
3751
[=]() {
3852
struct_acc.member_acc[2].use();
3953
});
54+
55+
a_kernel<class kernel_D>(
56+
[=]() {
57+
foo local = struct_array[1];
58+
});
4059
}
4160

4261
// Check kernel_A parameters
@@ -81,8 +100,8 @@ int main() {
81100
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
82101
// CHECK-NEXT: CompoundStmt
83102
// CHECK-NEXT: DeclStmt
84-
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp:37:7)' cinit
85-
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:37:7)'
103+
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
104+
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'
86105
// CHECK-NEXT: InitListExpr {{.*}} 'struct_acc_t'
87106
// CHECK-NEXT: InitListExpr {{.*}} 'Accessor [2]'
88107
// CHECK-NEXT: CXXConstructExpr {{.*}} 'Accessor [2]'
@@ -93,3 +112,92 @@ int main() {
93112
// CHECK-NEXT: MemberExpr {{.*}}__init
94113
// CHECK: CXXMemberCallExpr {{.*}} 'void'
95114
// CHECK-NEXT: MemberExpr {{.*}}__init
115+
116+
// Check kernel_D parameters
117+
// CHECK: FunctionDecl {{.*}}kernel_D{{.*}} 'void (int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int)'
118+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_a 'int'
119+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
120+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
121+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
122+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
123+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
124+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
125+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
126+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
127+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_c 'int'
128+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_a 'int'
129+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
130+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
131+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
132+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
133+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
134+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
135+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
136+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
137+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_c 'int'
138+
// CHECK-NEXT: CompoundStmt
139+
// CHECK-NEXT: DeclStmt
140+
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
141+
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'
142+
143+
// Initializer for struct array i.e. foo struct_array[2]
144+
// CHECK-NEXT: InitListExpr {{.*}} 'foo [2]'
145+
146+
// Initializer for first element of struct_array
147+
// CHECK-NEXT: InitListExpr {{.*}} 'foo'
148+
// CHECK-NEXT: ImplicitCastExpr
149+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_a' 'int'
150+
// Initializer for struct array inside foo i.e. foo_inner foo_b[2]
151+
// CHECK-NEXT: InitListExpr {{.*}} 'foo_inner [2]'
152+
// Initializer for first element of inner struct array
153+
// CHECK-NEXT: InitListExpr {{.*}} 'foo_inner'
154+
// CHECK-NEXT: ImplicitCastExpr
155+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int'
156+
// CHECK-NEXT: ImplicitCastExpr
157+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int'
158+
// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
159+
// CHECK-NEXT: ImplicitCastExpr
160+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
161+
// CHECK-NEXT: ImplicitCastExpr
162+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
163+
// Initializer for second element of inner struct array
164+
// CHECK-NEXT: InitListExpr {{.*}} 'foo_inner'
165+
// CHECK-NEXT: ImplicitCastExpr
166+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int'
167+
// CHECK-NEXT: ImplicitCastExpr
168+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int'
169+
// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
170+
// CHECK-NEXT: ImplicitCastExpr
171+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
172+
// CHECK-NEXT: ImplicitCastExpr
173+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
174+
// CHECK-NEXT: ImplicitCastExpr
175+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int'
176+
177+
// Initializer for second element of struct_array
178+
// CHECK-NEXT: InitListExpr {{.*}} 'foo'
179+
// CHECK-NEXT: ImplicitCastExpr
180+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_a' 'int'
181+
// CHECK-NEXT: InitListExpr {{.*}} 'foo_inner [2]'
182+
// CHECK-NEXT: InitListExpr {{.*}} 'foo_inner'
183+
// CHECK-NEXT: ImplicitCastExpr
184+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int'
185+
// CHECK-NEXT: ImplicitCastExpr
186+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int'
187+
// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
188+
// CHECK-NEXT: ImplicitCastExpr
189+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
190+
// CHECK-NEXT: ImplicitCastExpr
191+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
192+
// CHECK-NEXT: InitListExpr {{.*}} 'foo_inner'
193+
// CHECK-NEXT: ImplicitCastExpr
194+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int'
195+
// CHECK-NEXT: ImplicitCastExpr
196+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int'
197+
// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
198+
// CHECK-NEXT: ImplicitCastExpr
199+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
200+
// CHECK-NEXT: ImplicitCastExpr
201+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
202+
// CHECK-NEXT: ImplicitCastExpr
203+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int'

0 commit comments

Comments
 (0)