Skip to content

[SYCL] Fix crash when kernel object field is an array of structs #2083

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 6 commits into from
Jul 14, 2020
Merged

[SYCL] Fix crash when kernel object field is an array of structs #2083

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
c9e371d
Fix crash on array of structs.
elizabethandrews Jul 9, 2020
4b15731
Remove unnecessary option
elizabethandrews Jul 10, 2020
ac93254
Corrected code to work for array of streams.
elizabethandrews Jul 10, 2020
707df53
Remove extra lines
elizabethandrews Jul 10, 2020
95ba2fe
Use getBaseElementTypeUnsafe to replace unnecessary code
elizabethandrews Jul 10, 2020
88116ba
Add missing CHECK directives.
elizabethandrews Jul 13, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 4 additions & 4 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1584,9 +1584,6 @@ class SyclKernelBodyCreator
}

void addStructInit(const CXXRecordDecl *RD) {
if (!RD)
return;

const ASTRecordLayout &Info =
SemaRef.getASTContext().getASTRecordLayout(RD);
int NumberOfFields = Info.getFieldCount();
Expand All @@ -1607,7 +1604,10 @@ class SyclKernelBodyCreator
}

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

// Initializers for accessors inside stream not added.
if (!Util::isSyclStreamType(FD->getType()))
Expand Down
112 changes: 111 additions & 1 deletion clang/test/CodeGenSYCL/kernel-param-pod-array.cpp
100755 → 100644
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,14 +11,31 @@ __attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
kernelFunc();
}

struct foo_inner {
int foo_inner_x;
int foo_inner_y;
};

struct foo {
int foo_a;
foo_inner foo_b[2];
int foo_c;
};

int main() {

int a[2];
foo struct_array[2];

a_kernel<class kernel_B>(
[=]() {
int local = a[1];
});

a_kernel<class kernel_C>(
[=]() {
foo local = struct_array[1];
});
}

// Check kernel_B parameters
Expand All @@ -40,4 +57,97 @@ int main() {
// CHECK: store i32 [[ARRAY0]], i32* [[ARRAY_BEGIN]], align 4
// CHECK: [[ARRAY_ELEMENT:%[a-zA-Z_.]+]] = getelementptr inbounds i32, i32* %arrayinit.begin, i64 1
// CHECK: [[ARRAY1:%[0-9]*]] = load i32, i32* [[ELEM_L1]], align 4
// CHECK: store i32 [[ARRAY1]], i32* [[ARRAY_ELEMENT]], align 4
// CHECK: store i32 [[ARRAY1]], i32* [[ARRAY_ELEMENT]], align 4

// Check kernel_C parameters
// CHECK: define spir_kernel void @{{.*}}kernel_C
// 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_]+]],
// 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_]+]]

// Check local lambda object alloca
// CHECK: [[KERNEL_OBJ:%[0-9]+]] = alloca %"class.{{.*}}.anon.0", align 4

// Check local stores
// CHECK: store i32 [[FOO1_A]], i32* [[FOO1_A_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
// CHECK: store i32 [[FOO1_B1_X]], i32* [[FOO1_B1_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
// CHECK: store i32 [[FOO1_B1_Y]], i32* [[FOO1_B1_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
// CHECK: store i32 [[FOO1_B2_X]], i32* [[FOO1_B2_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
// CHECK: store i32 [[FOO1_B2_Y]], i32* [[FOO1_B2_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
// CHECK: store i32 [[FOO1_C]], i32* [[FOO1_C_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
// CHECK: store i32 [[FOO2_A]], i32* [[FOO2_A_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
// CHECK: store i32 [[FOO2_B1_X]], i32* [[FOO2_B1_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
// CHECK: store i32 [[FOO2_B1_Y]], i32* [[FOO2_B1_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
// CHECK: store i32 [[FOO2_B2_X]], i32* [[FOO2_B2_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
// CHECK: store i32 [[FOO2_B2_Y]], i32* [[FOO2_B2_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
// CHECK: store i32 [[FOO2_C]], i32* [[FOO2_C_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4

// Check initialization of local array

// Initialize struct_array[0].foo_a
// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon.0", %"class.{{.*}}.anon.0"* [[KERNEL_OBJ]], i32 0, i32 0
// CHECK: [[FOO_ARRAY_0:%[a-zA-Z_.]+]] = getelementptr inbounds [2 x %struct.{{.*}}.foo], [2 x %struct.{{.*}}.foo]* [[GEP]], i64 0, i64 0
// CHECK: [[GEP_FOO1_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* [[FOO_ARRAY_0]], i32 0, i32 0
// CHECK: [[LOAD_FOO1_A:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_A_LOCAL]], align 4
// CHECK: store i32 [[LOAD_FOO1_A]], i32* [[GEP_FOO1_A]], align 4

// Initialize struct_array[0].foo_b[0].x
// CHECK: [[GEP_FOO1_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* [[FOO_ARRAY_0]], i32 0, i32 1
// 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
// 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
// CHECK: [[LOAD_FOO1_B1_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B1_X_LOCAL]], align 4
// CHECK: store i32 [[LOAD_FOO1_B1_X]], i32* [[GEP_FOO1_B1_X]], align 4

// Initialize struct_array[0].foo_b[0].y
// 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
// CHECK: [[LOAD_FOO1_B1_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B1_Y_LOCAL]], align 4
// CHECK: store i32 [[LOAD_FOO1_B1_Y]], i32* [[GEP_FOO1_B1_Y]], align 4

// Initialize struct_array[0].foo_b[1].x
// 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
// 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
// CHECK: [[LOAD_FOO1_B2_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B2_X_LOCAL]], align 4
// CHECK: store i32 [[LOAD_FOO1_B2_X]], i32* [[GEP_FOO1_B2_X]], align 4

// Initialize struct_array[0].foo_b[1].y
// 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
// CHECK: [[LOAD_FOO1_B2_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B2_Y_LOCAL]], align 4
// CHECK: store i32 [[LOAD_FOO1_B2_Y]], i32* [[GEP_FOO1_B2_Y]], align 4

// Initialize struct_array[0].foo_c
// CHECK: [[GEP_FOO1_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_0]], i32 0, i32 2
// CHECK: [[LOAD_FOO1_C:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_C_LOCAL]], align 4
// CHECK: store i32 [[LOAD_FOO1_C]], i32* [[GEP_FOO1_C]], align 4

// Initialize struct_array[1].foo_a
// CHECK: [[FOO_ARRAY_1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct._ZTS3foo.foo, %struct._ZTS3foo.foo* [[FOO_ARRAY_0]], i64 1
// CHECK: [[GEP_FOO2_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_1]], i32 0, i32 0
// CHECK: [[LOAD_FOO2_A:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_A_LOCAL]], align 4
// CHECK: store i32 [[LOAD_FOO2_A]], i32* [[GEP_FOO2_A]], align 4

// Initialize struct_array[1].foo_b[0].x
// CHECK: [[GEP_FOO2_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* [[FOO_ARRAY_1]], i32 0, i32 1
// 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
// 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
// CHECK: [[LOAD_FOO2_B1_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B1_X_LOCAL]], align 4
// CHECK: store i32 [[LOAD_FOO2_B1_X]], i32* [[GEP_FOO2_B1_X]]

// Initialize struct_array[1].foo_b[0].y
// 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
// CHECK: [[LOAD_FOO2_B1_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B1_Y_LOCAL]], align 4
// CHECK: store i32 [[LOAD_FOO2_B1_Y]], i32* [[GEP_FOO2_B1_Y]], align 4

// Initialize struct_array[1].foo_b[1].x
// 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
// 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
// CHECK: [[LOAD_FOO2_B2_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B2_X_LOCAL]], align 4
// CHECK: store i32 [[LOAD_FOO2_B2_X]], i32* [[GEP_FOO2_B2_X]], align 4

// Initialize struct_array[1].foo_b[1].y
// 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
// CHECK: [[LOAD_FOO2_B2_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B2_Y_LOCAL]], align 4
// CHECK: store i32 [[LOAD_FOO2_B2_Y]], i32* [[GEP_FOO2_B2_Y]], align 4

// Initialize struct_array[1].foo_c
// CHECK: [[GEP_FOO2_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_1]], i32 0, i32 2
// CHECK: [[LOAD_FOO2_C:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_C_LOCAL]], align 4
// CHECK: store i32 [[LOAD_FOO2_C]], i32* [[GEP_FOO2_C]], align 4
112 changes: 110 additions & 2 deletions clang/test/SemaSYCL/array-kernel-param.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,20 @@ int main() {
Accessor member_acc[2];
} struct_acc;

struct foo_inner {
int foo_inner_x;
int foo_inner_y;
int foo_inner_z[2];
};

struct foo {
int foo_a;
foo_inner foo_b[2];
int foo_c;
};

foo struct_array[2];

a_kernel<class kernel_A>(
[=]() {
acc[1].use();
Expand All @@ -37,6 +51,11 @@ int main() {
[=]() {
struct_acc.member_acc[2].use();
});

a_kernel<class kernel_D>(
[=]() {
foo local = struct_array[1];
});
}

// Check kernel_A parameters
Expand Down Expand Up @@ -81,8 +100,8 @@ int main() {
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
// CHECK-NEXT: CompoundStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp:37:7)' cinit
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:37:7)'
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'
// CHECK-NEXT: InitListExpr {{.*}} 'struct_acc_t'
// CHECK-NEXT: InitListExpr {{.*}} 'Accessor [2]'
// CHECK-NEXT: CXXConstructExpr {{.*}} 'Accessor [2]'
Expand All @@ -93,3 +112,92 @@ int main() {
// CHECK-NEXT: MemberExpr {{.*}}__init
// CHECK: CXXMemberCallExpr {{.*}} 'void'
// CHECK-NEXT: MemberExpr {{.*}}__init

// Check kernel_D parameters
// 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)'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_a 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_c 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_a 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_c 'int'
// CHECK-NEXT: CompoundStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'

// Initializer for struct array i.e. foo struct_array[2]
// CHECK-NEXT: InitListExpr {{.*}} 'foo [2]'

// Initializer for first element of struct_array
// CHECK-NEXT: InitListExpr {{.*}} 'foo'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_a' 'int'
// Initializer for struct array inside foo i.e. foo_inner foo_b[2]
// CHECK-NEXT: InitListExpr {{.*}} 'foo_inner [2]'
// Initializer for first element of inner struct array
// CHECK-NEXT: InitListExpr {{.*}} 'foo_inner'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int'
// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
// Initializer for second element of inner struct array
// CHECK-NEXT: InitListExpr {{.*}} 'foo_inner'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int'
// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int'

// Initializer for second element of struct_array
// CHECK-NEXT: InitListExpr {{.*}} 'foo'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_a' 'int'
// CHECK-NEXT: InitListExpr {{.*}} 'foo_inner [2]'
// CHECK-NEXT: InitListExpr {{.*}} 'foo_inner'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int'
// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
// CHECK-NEXT: InitListExpr {{.*}} 'foo_inner'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int'
// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
// CHECK-NEXT: ImplicitCastExpr
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int'