Skip to content

Commit 5465fc0

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web' (#8)
2 parents 4bddbdb + b00fb7c commit 5465fc0

File tree

27 files changed

+727
-87
lines changed

27 files changed

+727
-87
lines changed

‎clang/lib/CodeGen/CodeGenModule.cpp

Lines changed: 7 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -2634,11 +2634,6 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
26342634
}
26352635
}
26362636

2637-
if (LangOpts.SYCLIsDevice && MustBeEmitted(Global)) {
2638-
addDeferredDeclToEmit(GD);
2639-
return;
2640-
}
2641-
26422637
// Ignore declarations, they will be emitted on their first use.
26432638
if (const auto *FD = dyn_cast<FunctionDecl>(Global)) {
26442639
// Forward declarations are emitted lazily on first use.
@@ -2690,6 +2685,13 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
26902685
}
26912686
}
26922687

2688+
// clang::ParseAST ensures that we emit the SYCL devices at the end, so
2689+
// anything that is a device (or indirectly called) will be handled later.
2690+
if (LangOpts.SYCLIsDevice && MustBeEmitted(Global)) {
2691+
addDeferredDeclToEmit(GD);
2692+
return;
2693+
}
2694+
26932695
// Defer code generation to first use when possible, e.g. if this is an inline
26942696
// function. If the global must always be emitted, do it eagerly if possible
26952697
// to benefit from cache locality.

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 9 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1244,7 +1244,11 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
12441244
// space, because OpenCL requires it.
12451245
QualType PointeeTy = FieldTy->getPointeeType();
12461246
Qualifiers Quals = PointeeTy.getQualifiers();
1247-
Quals.setAddressSpace(LangAS::opencl_global);
1247+
auto AS = Quals.getAddressSpace();
1248+
// Leave global_device and global_host address spaces as is to help FPGA
1249+
// device in memory allocations
1250+
if (AS != LangAS::opencl_global_device && AS != LangAS::opencl_global_host)
1251+
Quals.setAddressSpace(LangAS::opencl_global);
12481252
PointeeTy = SemaRef.getASTContext().getQualifiedType(
12491253
PointeeTy.getUnqualifiedType(), Quals);
12501254
QualType ModTy = SemaRef.getASTContext().getPointerType(PointeeTy);
@@ -1594,9 +1598,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
15941598
}
15951599

15961600
void addStructInit(const CXXRecordDecl *RD) {
1597-
if (!RD)
1598-
return;
1599-
16001601
const ASTRecordLayout &Info =
16011602
SemaRef.getASTContext().getASTRecordLayout(RD);
16021603
int NumberOfFields = Info.getFieldCount();
@@ -1617,7 +1618,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
16171618
}
16181619

16191620
bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD) final {
1620-
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();
16211625

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

‎clang/test/CodeGenSYCL/emit-all-decls-crash.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,6 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -femit-all-decls -o - | FileCheck %s
2+
3+
// This should not crash and we should not emit this declaration, even though
4+
// we have 'emit-all-decls'.
5+
// CHECK-NOT: define
6+
void foo(void);

‎clang/test/CodeGenSYCL/kernel-device-space-arg.cpp

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,16 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -disable-llvm-passes -o - | FileCheck %s
2+
3+
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function(i32 addrspace(5)* {{.*}} i32 addrspace(6)* {{.*}}
4+
5+
#include "sycl.hpp"
6+
7+
int main() {
8+
__attribute__((opencl_global_device)) int *GLOBDEV = nullptr;
9+
__attribute__((opencl_global_host)) int *GLOBHOST = nullptr;
10+
cl::sycl::kernel_single_task<class kernel_function>(
11+
[=]() {
12+
__attribute__((opencl_global_device)) int *DevPtr = GLOBDEV;
13+
__attribute__((opencl_global_host)) int *HostPtr = GLOBHOST;
14+
});
15+
return 0;
16+
}

‎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'

‎libclc/cmake/modules/AddLibclc.cmake

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -213,6 +213,8 @@ function(add_libclc_sycl_binding OUT_LIST)
213213
if( EXISTS ${SYCLDEVICE_BINDING} )
214214
set( SYCLDEVICE_BINDING_OUT ${CMAKE_CURRENT_BINARY_DIR}/sycldevice-binding-${ARG_TRIPLE}/sycldevice-binding.bc )
215215
add_custom_command( OUTPUT ${SYCLDEVICE_BINDING_OUT}
216+
COMMAND ${CMAKE_COMMAND} -E make_directory
217+
${CMAKE_CURRENT_BINARY_DIR}/sycldevice-binding-${ARG_TRIPLE}
216218
COMMAND ${LLVM_CLANG}
217219
-target ${ARG_TRIPLE}-sycldevice
218220
-fsycl

‎libdevice/fallback-cmath-fp64.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,9 @@
1010

1111
#ifdef __SPIR__
1212

13+
// To support fallback device libraries on-demand loading, please update the
14+
// DeviceLibFuncMap in llvm/tools/sycl-post-link/sycl-post-link.cpp if you add
15+
// or remove any item in this file.
1316
DEVICE_EXTERN_C
1417
double __devicelib_log(double x) { return __spirv_ocl_log(x); }
1518

‎libdevice/fallback-cmath.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,11 @@
1010

1111
#ifdef __SPIR__
1212

13+
// To support fallback device libraries on-demand loading, please update the
14+
// DeviceLibFuncMap in llvm/tools/sycl-post-link/sycl-post-link.cpp if you add
15+
// or remove any item in this file.
16+
// TODO: generate the DeviceLibFuncMap in sycl-post-link.cpp automatically
17+
// during the build based on libdevice to avoid manually sync.
1318
DEVICE_EXTERN_C
1419
float __devicelib_scalbnf(float x, int n) { return __spirv_ocl_ldexp(x, n); }
1520

‎libdevice/fallback-complex-fp64.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,9 @@
1212
#ifdef __SPIR__
1313
#include <cmath>
1414

15+
// To support fallback device libraries on-demand loading, please update the
16+
// DeviceLibFuncMap in llvm/tools/sycl-post-link/sycl-post-link.cpp if you add
17+
// or remove any item in this file.
1518
DEVICE_EXTERN_C
1619
double __devicelib_creal(double __complex__ z) { return __real__(z); }
1720

0 commit comments

Comments
 (0)