Skip to content

Commit 494c1fb

Browse files
elizabethandrewselizabethandrews
committed
[SYCL] Fix crash when kernel argument is a multi-dimensional array.
This patch fixes crash due to incorrect InitializedEntity for multi-dimensional arrays. When generating the InitializedEntity for an element, it is necessary to descend the array. For example, the initialized entity for s.array[x][y][z] is constructed using initialized entities for s.array[x][y], s.array[x] and s.array. Prior to this patch, the 'descending' was not done. Patch by: Rajiv Deodhar and Elizabeth Andrews Signed-off-by: Elizabeth Andrews <[email protected]>
1 parent 3084982 commit 494c1fb

File tree

3 files changed

+145
-19
lines changed

3 files changed

+145
-19
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 57 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1739,17 +1739,55 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
17391739
InitExprs.push_back(MemberInit.get());
17401740
}
17411741

1742+
int getDims() {
1743+
int Dims = 0;
1744+
for (int i = MemberExprBases.size() - 1; i >= 0; --i) {
1745+
if (!isa<ArraySubscriptExpr>(MemberExprBases[i]))
1746+
break;
1747+
++Dims;
1748+
}
1749+
return Dims;
1750+
}
1751+
1752+
int64_t getArrayIndex(int i) {
1753+
ArraySubscriptExpr *LastArrayRef =
1754+
cast<ArraySubscriptExpr>(MemberExprBases[i]);
1755+
Expr *LastIdx = LastArrayRef->getIdx();
1756+
llvm::APSInt Result;
1757+
SemaRef.VerifyIntegerConstantExpression(LastIdx, &Result);
1758+
return Result.getExtValue();
1759+
}
1760+
17421761
void createExprForScalarElement(FieldDecl *FD) {
1743-
InitializedEntity ArrayEntity =
1762+
llvm::SmallVector<InitializedEntity, 4> InitEntities;
1763+
1764+
InitializedEntity MemberEntity =
17441765
InitializedEntity::InitializeMember(FD, &VarEntity);
1766+
InitializedEntity Entity = InitializedEntity::InitializeElement(
1767+
SemaRef.getASTContext(), ArrayIndex, MemberEntity);
1768+
InitEntities.push_back(Entity);
1769+
// For multi-dimensional arrays, an initialized entity needs to be
1770+
// generated for each 'dimension'. For example, the initialized entity
1771+
// for s.array[x][y][z] is constructed using initialized entities for
1772+
// s.array[x][y], s.array[x] and s.array. InitEntities is used to maintain
1773+
// this. MemberExprBases is used to get the array index for 'current
1774+
// dimension'.
1775+
for (int i = MemberExprBases.size() - 2; i >= 0; --i) {
1776+
if (!isa<ArraySubscriptExpr>(MemberExprBases[i]))
1777+
break;
1778+
InitializedEntity NewEntity = InitializedEntity::InitializeElement(
1779+
SemaRef.getASTContext(), getArrayIndex(i), InitEntities.back());
1780+
InitEntities.push_back(NewEntity);
1781+
}
1782+
1783+
ArrayIndex++;
1784+
17451785
InitializationKind InitKind =
17461786
InitializationKind::CreateCopy(SourceLocation(), SourceLocation());
17471787
Expr *DRE = createInitExpr(FD);
1748-
InitializedEntity Entity = InitializedEntity::InitializeElement(
1749-
SemaRef.getASTContext(), ArrayIndex, ArrayEntity);
1750-
ArrayIndex++;
1751-
InitializationSequence InitSeq(SemaRef, Entity, InitKind, DRE);
1752-
ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, DRE);
1788+
InitializationSequence InitSeq(SemaRef, InitEntities.back(), InitKind, DRE);
1789+
ExprResult MemberInit =
1790+
InitSeq.Perform(SemaRef, InitEntities.back(), InitKind, DRE);
17531791
InitExprs.push_back(MemberInit.get());
17541792
}
17551793

@@ -1763,7 +1801,14 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
17631801
Expr *ILE = new (SemaRef.getASTContext())
17641802
InitListExpr(SemaRef.getASTContext(), SourceLocation(), ArrayInitExprs,
17651803
SourceLocation());
1766-
ILE->setType(FD->getType());
1804+
QualType ILEType = FD->getType();
1805+
for (int i = getDims(); i > 1; i--) {
1806+
const ConstantArrayType *CAT =
1807+
SemaRef.getASTContext().getAsConstantArrayType(ILEType);
1808+
assert(CAT && "Should only be called on constant-size array.");
1809+
ILEType = CAT->getElementType();
1810+
}
1811+
ILE->setType(ILEType);
17671812
InitExprs.push_back(ILE);
17681813
}
17691814

@@ -2027,15 +2072,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
20272072
}
20282073

20292074
bool nextElement(QualType ET) final {
2030-
ArraySubscriptExpr *LastArrayRef =
2031-
cast<ArraySubscriptExpr>(MemberExprBases.back());
2075+
// Top of MemberExprBases holds ArraySubscriptExpression of element
2076+
// we just finished processing.
2077+
int64_t nextIndex = getArrayIndex((MemberExprBases.size() - 1)) + 1;
20322078
MemberExprBases.pop_back();
2033-
Expr *LastIdx = LastArrayRef->getIdx();
2034-
llvm::APSInt Result;
2035-
SemaRef.VerifyIntegerConstantExpression(LastIdx, &Result);
20362079
Expr *ArrayBase = MemberExprBases.back();
2037-
ExprResult IndexExpr = SemaRef.ActOnIntegerConstant(
2038-
SourceLocation(), Result.getExtValue() + 1);
2080+
ExprResult IndexExpr =
2081+
SemaRef.ActOnIntegerConstant(SourceLocation(), nextIndex);
20392082
ExprResult ElementBase = SemaRef.CreateBuiltinArraySubscriptExpr(
20402083
ArrayBase, SourceLocation(), IndexExpr.get(), SourceLocation());
20412084
MemberExprBases.push_back(ElementBase.get());

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

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,7 @@ struct foo {
2525
int main() {
2626

2727
int a[2];
28+
int array_2D[2][1];
2829
foo struct_array[2];
2930

3031
a_kernel<class kernel_B>(
@@ -36,6 +37,11 @@ int main() {
3637
[=]() {
3738
foo local = struct_array[1];
3839
});
40+
41+
a_kernel<class kernel_D>(
42+
[=]() {
43+
int local = array_2D[0][0];
44+
});
3945
}
4046

4147
// Check kernel_B parameters
@@ -151,3 +157,25 @@ int main() {
151157
// CHECK: [[GEP_FOO2_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_1]], i32 0, i32 2
152158
// CHECK: [[LOAD_FOO2_C:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_C_LOCAL]], align 4
153159
// CHECK: store i32 [[LOAD_FOO2_C]], i32* [[GEP_FOO2_C]], align 4
160+
161+
// Check kernel_D parameters
162+
// CHECK: define spir_kernel void @{{.*}}kernel_D
163+
// CHECK-SAME: i32 [[ARR_2D_1:%[a-zA-Z0-9_]+]], i32 [[ARR_2D_2:%[a-zA-Z0-9_]+]]
164+
165+
// Check local lambda object alloca
166+
// CHECK: [[LAMBDA_OBJ:%[0-9]+]] = alloca %"class.{{.*}}.anon.1", align 4
167+
168+
// Check local stores
169+
// CHECK: store i32 [[ARR_2D_1]], i32* [[ARR_2D_1_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
170+
// CHECK: store i32 [[ARR_2D_2]], i32* [[ARR_2D_2_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4
171+
172+
// Check initialization of local array
173+
// CHECK: [[GEP_ARR_2D:%[0-9]*]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon.1", %"class._ZTSZ4mainE3$_0.anon.1"* [[LAMBDA_OBJ]], i32 0, i32 0
174+
// CHECK: [[GEP_ARR_BEGIN1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [2 x [1 x i32]], [2 x [1 x i32]]* [[GEP_ARR_2D]], i64 0, i64 0
175+
// CHECK: [[GEP_ARR_ELEM0:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [1 x i32], [1 x i32]* [[GEP_ARR_BEGIN1]], i64 0, i64 0
176+
// CHECK: [[ARR_2D_ELEM0:%[0-9]*]] = load i32, i32* [[ARR_2D_1_LOCAL]], align 4
177+
// CHECK: store i32 [[ARR_2D_ELEM0]], i32* [[GEP_ARR_ELEM0]], align 4
178+
// CHECK: [[GEP_ARR_BEGIN2:%[a-zA-Z_.]+]] = getelementptr inbounds [1 x i32], [1 x i32]* [[GEP_ARR_BEGIN1]], i64 1
179+
// CHECK: [[GEP_ARR_ELEM1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [1 x i32], [1 x i32]* [[GEP_ARR_BEGIN2]], i64 0, i64 0
180+
// CHECK: [[ARR_2D_ELEM1:%[0-9]*]] = load i32, i32* [[ARR_2D_2_LOCAL]], align 4
181+
// CHECK: store i32 [[ARR_2D_ELEM1]], i32* [[GEP_ARR_ELEM1]], align 4

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

Lines changed: 60 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -38,11 +38,14 @@ int main() {
3838
struct foo {
3939
int foo_a;
4040
foo_inner foo_b[2];
41+
int foo_2D[2][1];
4142
int foo_c;
4243
};
4344

4445
foo struct_array[2];
4546

47+
int array_2D[2][3];
48+
4649
a_kernel<class kernel_A>(
4750
[=]() {
4851
acc[1].use();
@@ -67,6 +70,11 @@ int main() {
6770
[=]() {
6871
int local = s.a[2];
6972
});
73+
74+
a_kernel<class kernel_F>(
75+
[=]() {
76+
int local = array_2D[1][1];
77+
});
7078
}
7179

7280
// Check kernel_A parameters
@@ -111,8 +119,8 @@ int main() {
111119
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
112120
// CHECK-NEXT: CompoundStmt
113121
// CHECK-NEXT: DeclStmt
114-
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp:57:7)' cinit
115-
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:57:7)'
122+
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
123+
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'
116124
// CHECK-NEXT: InitListExpr {{.*}} 'struct_acc_t'
117125
// CHECK-NEXT: InitListExpr {{.*}} 'Accessor [2]'
118126
// CHECK-NEXT: CXXConstructExpr {{.*}} 'Accessor [2]'
@@ -125,7 +133,7 @@ int main() {
125133
// CHECK-NEXT: MemberExpr {{.*}}__init
126134

127135
// Check kernel_D parameters
128-
// 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)'
136+
// 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, int, int, int, int)'
129137
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_a 'int'
130138
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
131139
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
@@ -135,6 +143,8 @@ int main() {
135143
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
136144
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
137145
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
146+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D 'int'
147+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D 'int'
138148
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_c 'int'
139149
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_a 'int'
140150
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int'
@@ -145,6 +155,8 @@ int main() {
145155
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int'
146156
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
147157
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int'
158+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D 'int'
159+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D 'int'
148160
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_c 'int'
149161
// CHECK-NEXT: CompoundStmt
150162
// CHECK-NEXT: DeclStmt
@@ -182,6 +194,13 @@ int main() {
182194
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
183195
// CHECK-NEXT: ImplicitCastExpr
184196
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
197+
// CHECK-NEXT: InitListExpr {{.*}} 'int [2][1]'
198+
// CHECK-NEXT: InitListExpr {{.*}} 'int [1]'
199+
// CHECK-NEXT: ImplicitCastExpr
200+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_2D' 'int'
201+
// CHECK-NEXT: InitListExpr {{.*}} 'int [1]'
202+
// CHECK-NEXT: ImplicitCastExpr
203+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_2D' 'int'
185204
// CHECK-NEXT: ImplicitCastExpr
186205
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int'
187206

@@ -210,6 +229,13 @@ int main() {
210229
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
211230
// CHECK-NEXT: ImplicitCastExpr
212231
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int'
232+
// CHECK-NEXT: InitListExpr {{.*}} 'int [2][1]'
233+
// CHECK-NEXT: InitListExpr {{.*}} 'int [1]'
234+
// CHECK-NEXT: ImplicitCastExpr
235+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_2D' 'int'
236+
// CHECK-NEXT: InitListExpr {{.*}} 'int [1]'
237+
// CHECK-NEXT: ImplicitCastExpr
238+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_2D' 'int'
213239
// CHECK-NEXT: ImplicitCastExpr
214240
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int'
215241

@@ -220,8 +246,8 @@ int main() {
220246
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_a 'int':'int'
221247
// CHECK-NEXT: CompoundStmt
222248
// CHECK-NEXT: DeclStmt
223-
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp:67:7)' cinit
224-
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:67:7)'
249+
// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit
250+
// CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})'
225251
// CHECK-NEXT: InitListExpr {{.*}} 'S<int>'
226252
// CHECK-NEXT: InitListExpr {{.*}} 'int [3]'
227253
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int':'int'
@@ -230,3 +256,32 @@ int main() {
230256
// CHECK-NEXT: DeclRefExpr {{.*}} 'int':'int'
231257
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int':'int'
232258
// CHECK-NEXT: DeclRefExpr {{.*}} 'int':'int'
259+
260+
// Check kernel_F parameters
261+
// CHECK: FunctionDecl {{.*}}kernel_F{{.*}} 'void (int, int, int, int, int, int)'
262+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
263+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
264+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
265+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
266+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
267+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
268+
// Check kernel_F inits
269+
// CHECK-NEXT: CompoundStmt
270+
// CHECK-NEXT: DeclStmt
271+
// CHECK-NEXT: VarDecl {{.*}} cinit
272+
// CHECK-NEXT: InitListExpr
273+
// CHECK-NEXT: InitListExpr {{.*}} 'int [2][3]'
274+
// CHECK-NEXT: InitListExpr {{.*}} 'int [3]'
275+
// CHECK-NEXT: ImplicitCastExpr
276+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
277+
// CHECK-NEXT: ImplicitCastExpr
278+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
279+
// CHECK-NEXT: ImplicitCastExpr
280+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
281+
// CHECK-NEXT: InitListExpr {{.*}} 'int [3]'
282+
// CHECK-NEXT: ImplicitCastExpr
283+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
284+
// CHECK-NEXT: ImplicitCastExpr
285+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
286+
// CHECK-NEXT: ImplicitCastExpr
287+
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'

0 commit comments

Comments
 (0)