Skip to content

Commit 52ce3f2

Browse files
rdeodharrdeodhar
committed
Updated support for arrays.
1 parent 546c58d commit 52ce3f2

File tree

6 files changed

+493
-81
lines changed

6 files changed

+493
-81
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 23 additions & 43 deletions
Original file line numberDiff line numberDiff line change
@@ -644,6 +644,22 @@ static ParamDesc makeParamDesc(ASTContext &Ctx, const CXXBaseSpecifier &Src,
644644
Ctx.getTrivialTypeSourceInfo(Ty));
645645
}
646646

647+
// Create a new class around a field - used to wrap arrays.
648+
static RecordDecl *wrapAnArray(ASTContext &Ctx, const QualType ArgTy,
649+
FieldDecl *&Field) {
650+
RecordDecl *NewClass = Ctx.buildImplicitRecord("wrapped_array");
651+
NewClass->startDefinition();
652+
Field = FieldDecl::Create(
653+
Ctx, NewClass, SourceLocation(), SourceLocation(),
654+
/*Id=*/nullptr, ArgTy,
655+
Ctx.getTrivialTypeSourceInfo(ArgTy, SourceLocation()),
656+
/*BW=*/nullptr, /*Mutable=*/false, /*InitStyle=*/ICIS_NoInit);
657+
Field->setAccess(AS_public);
658+
NewClass->addDecl(Field);
659+
NewClass->completeDefinition();
660+
return NewClass;
661+
}
662+
647663
/// \return the target of given SYCL accessor type
648664
static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) {
649665
return static_cast<target>(
@@ -1035,23 +1051,6 @@ class SyclKernelDeclCreator
10351051
return true;
10361052
}
10371053

1038-
// Create a new class around a field - used to wrap arrays.
1039-
RecordDecl *wrapAnArray(const QualType ArgTy, FieldDecl *Field) {
1040-
RecordDecl *NewClass =
1041-
SemaRef.getASTContext().buildImplicitRecord("wrapped_array");
1042-
NewClass->startDefinition();
1043-
Field = FieldDecl::Create(
1044-
SemaRef.getASTContext(), NewClass, SourceLocation(), SourceLocation(),
1045-
/*Id=*/nullptr, ArgTy,
1046-
SemaRef.getASTContext().getTrivialTypeSourceInfo(ArgTy,
1047-
SourceLocation()),
1048-
/*BW=*/nullptr, /*Mutable=*/false, /*InitStyle=*/ICIS_NoInit);
1049-
Field->setAccess(AS_public);
1050-
NewClass->addDecl(Field);
1051-
NewClass->completeDefinition();
1052-
return NewClass;
1053-
};
1054-
10551054
static void setKernelImplicitAttrs(ASTContext &Context, FunctionDecl *FD,
10561055
StringRef Name) {
10571056
// Set implicit attributes.
@@ -1142,13 +1141,9 @@ class SyclKernelDeclCreator
11421141
}
11431142

11441143
bool handleArrayType(FieldDecl *FD, QualType FieldTy) final {
1145-
if (!cast<ConstantArrayType>(FieldTy)
1146-
->getElementType()
1147-
->isStructureOrClassType()) {
1148-
RecordDecl *NewClass = wrapAnArray(FieldTy, FD);
1149-
QualType ST = SemaRef.getASTContext().getRecordType(NewClass);
1150-
addParam(FD, ST);
1151-
}
1144+
RecordDecl *NewClass = wrapAnArray(SemaRef.getASTContext(), FieldTy, FD);
1145+
QualType ST = SemaRef.getASTContext().getRecordType(NewClass);
1146+
addParam(FD, ST);
11521147
return true;
11531148
}
11541149

@@ -1295,7 +1290,7 @@ class SyclKernelBodyCreator
12951290
// The first and only field of the wrapper struct is the array
12961291
FieldDecl *Array = *(WrapperStruct->field_begin());
12971292
Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue,
1298-
SourceLocation());
1293+
SourceLocation());
12991294
Expr *InitExpr = BuildMemberExpr(DRE, Array);
13001295
InitializationKind InitKind = InitializationKind::CreateDirect(
13011296
SourceLocation(), SourceLocation(), SourceLocation());
@@ -1307,15 +1302,6 @@ class SyclKernelBodyCreator
13071302
InitExprs.push_back(MemberInit.get());
13081303
}
13091304

1310-
void createExprForArrayElement(size_t ArrayIndex) {
1311-
Expr *ArrayBase = MemberExprBases.back();
1312-
ExprResult IndexExpr =
1313-
SemaRef.ActOnIntegerConstant(SourceLocation(), ArrayIndex);
1314-
ExprResult ElementBase = SemaRef.CreateBuiltinArraySubscriptExpr(
1315-
ArrayBase, SourceLocation(), IndexExpr.get(), SourceLocation());
1316-
MemberExprBases.push_back(ElementBase.get());
1317-
}
1318-
13191305
void createSpecialMethodCall(const CXXRecordDecl *SpecialClass, Expr *Base,
13201306
const std::string &MethodName,
13211307
FieldDecl *Field) {
@@ -1455,11 +1441,7 @@ class SyclKernelBodyCreator
14551441
}
14561442

14571443
bool handleArrayType(FieldDecl *FD, QualType FieldTy) final {
1458-
if (!cast<ConstantArrayType>(FieldTy)
1459-
->getElementType()
1460-
->isStructureOrClassType()) {
1461-
createExprForArray(FD);
1462-
}
1444+
createExprForArray(FD);
14631445
return true;
14641446
}
14651447

@@ -1600,10 +1582,8 @@ class SyclKernelIntHeaderCreator
16001582
}
16011583

16021584
bool handleArrayType(FieldDecl *FD, QualType FieldTy) final {
1603-
if (!cast<ConstantArrayType>(FieldTy)
1604-
->getElementType()
1605-
->isStructureOrClassType())
1606-
addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout);
1585+
wrapAnArray(SemaRef.getASTContext(), FieldTy, FD);
1586+
addParam(FD, FD->getType(), SYCLIntegrationHeader::kind_std_layout);
16071587
return true;
16081588
}
16091589

‎clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang -I %S/Inputs -fsycl-device-only -Xclang -fsycl-int-header=%t.h %s -c -o %T/kernel.spv
1+
// RUN: %clang -I %S/Inputs -fsycl-device-only -Xclang -fsycl-int-header=/iusers/rdeodhar/work/dpcc/jira/7004/t.h %s -c -o %T/kernel.spv
22
// RUN: FileCheck -input-file=%t.h %s
33

44
// This test checks the integration header generated when
@@ -19,7 +19,8 @@
1919

2020
// CHECK: static constexpr
2121
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
22-
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_A
22+
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_A
23+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 24, 0 },
2324
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 },
2425
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 },
2526
// CHECK-EMPTY:

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

Lines changed: 36 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -25,50 +25,55 @@ int main() {
2525

2626
// Check kernel_A parameters
2727
// CHECK: define spir_kernel void @{{.*}}kernel_A
28-
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]],
29-
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]],
30-
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]],
31-
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]],
32-
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG2:%[a-zA-Z0-9_]+_4]],
33-
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+_6]],
34-
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+_7]],
35-
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+_8]])
36-
37-
// Check alloca for pointer arguments
38-
// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8
39-
// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8
40-
41-
// Check lambda object alloca
28+
// CHECK-SAME: %struct.{{.*}}.wrapped_array* byval{{.*}}align 4 %_arg_,
29+
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+_1]],
30+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_2]],
31+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_3]],
32+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_4]],
33+
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG2:%[a-zA-Z0-9_]+_5]],
34+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+_7]],
35+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+_8]],
36+
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+_9]])
37+
38+
// CHECK alloca for pointer arguments
39+
// CHECK: [[MEM_ARG1:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8
40+
// CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8
41+
42+
// CHECK lambda object alloca
4243
// CHECK: [[LOCAL_OBJECT:%0]] = alloca %"class.{{.*}}.anon", align 4
4344

44-
// Check allocas for ranges
45+
// CHECK allocas for ranges
4546
// CHECK: [[ACC_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
4647
// CHECK: [[MEM_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
4748
// CHECK: [[OFFSET1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
4849
// CHECK: [[ACC_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
4950
// CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
5051
// CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
5152

52-
// Check accessor array GEP for acc[0]
53-
// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
54-
// CHECK: [[Z0:%[a-zA-Z0-9_]*]] = getelementptr inbounds{{.*}}[[ACCESSOR_ARRAY1]], i64 0, i64 0
53+
// Check array initialization
54+
// CHECK: arrayinit.body:
55+
// CHECK: arrayinit.end:
5556

56-
// Check load from kernel pointer argument alloca
57-
// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}
57+
// CHECK accessor array GEP for acc[0]
58+
// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
59+
// CHECK: [[Z0:%[a-zA-Z0-9_]*]] = getelementptr inbounds{{.*}}[[ACCESSOR_ARRAY1]], i64 0, i64 0
5860

59-
// Check acc[0] __init method call
60-
// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z0]] to %"class{{.*}}accessor" addrspace(4)*
61+
// CHECK load from kernel pointer argument alloca
62+
// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]]
6163

62-
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
64+
// CHECK acc[0] __init method call
65+
// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z0]] to %"class{{.*}}accessor" addrspace(4)*
6366

64-
// Check accessor array GEP for acc[1]
65-
// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
66-
// CHECK: [[Z1:%[a-zA-Z0-9_]*]] = getelementptr inbounds{{.*}}[[ACCESSOR_ARRAY2]], i64 0, i64 1
67+
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]])
6768

68-
// Check load from kernel pointer argument alloca
69-
// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}
69+
// CHECK accessor array GEP for acc[1]
70+
// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
71+
// CHECK: [[Z1:%[a-zA-Z0-9_]*]] = getelementptr inbounds{{.*}}[[ACCESSOR_ARRAY2]], i64 0, i64 1
7072

71-
// Check acc[1] __init method call
72-
// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z1]] to %"class{{.*}}accessor" addrspace(4)*
73+
// CHECK load from kernel pointer argument alloca
74+
// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG2]]
7375

74-
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]])
76+
// CHECK acc[1] __init method call
77+
// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z1]] to %"class{{.*}}accessor" addrspace(4)*
78+
79+
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]])

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

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -40,7 +40,8 @@ int main() {
4040
}
4141

4242
// Check kernel_A parameters
43-
// CHECK: FunctionDecl {{.*}}kernel_A{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
43+
// CHECK: FunctionDecl {{.*}}kernel_A{{.*}} 'void (wrapped_array, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
44+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'wrapped_array'
4445
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__global int *'
4546
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'cl::sycl::range<1>'
4647
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'cl::sycl::range<1>'
@@ -54,7 +55,7 @@ int main() {
5455
// CHECK: CXXMemberCallExpr {{.*}} 'void'
5556
// CHECK-NEXT: MemberExpr {{.*}}__init
5657

57-
// Check kernel_B parameters
58+
// CHECK kernel_B parameters
5859
// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (wrapped_array)'
5960
// CHECK-NEXT: ParmVarDecl {{.*}} 'wrapped_array'
6061
// CHECK-NEXT: CompoundStmt
@@ -63,7 +64,7 @@ int main() {
6364
// CHECK-NEXT: InitListExpr
6465
// CHECK-NEXT: ArrayInitLoopExpr {{.*}} 'int [100]'
6566

66-
// Check kernel_C parameters
67+
// CHECK kernel_C parameters
6768
// CHECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (struct {{.*}}, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
6869
// CHECK-NEXT: ParmVarDecl {{.*}} 'struct {{.*}}'
6970
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
@@ -83,7 +84,7 @@ int main() {
8384
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
8485
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
8586

86-
// Check that four accessor init functions are called
87+
// CHECK that four accessor init functions are called
8788
// CHECK: CXXMemberCallExpr {{.*}} 'void'
8889
// CHECK-NEXT: MemberExpr {{.*}}__init
8990
// CHECK: CXXMemberCallExpr {{.*}} 'void'

‎sycl/doc/CompilerAndRuntimeDesign.md

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -168,7 +168,8 @@ __kernel KernelName(global int* a) {
168168
```
169169

170170
OpenCL kernel function is generated by the compiler inside the Sema using AST
171-
nodes.
171+
nodes. Additional details of kernel parameter passing may be found in the document
172+
[SYCL Kernel Parameter Handling and Array Support](KernelParameterPassing.md) .
172173

173174
### SYCL support in the driver
174175

0 commit comments

Comments
 (0)