Skip to content

Commit 0b2de9e

Browse files
elizabethandrewselizabethandrews
authored
[SYCL] Decompose kernel parameters and add inheritance support (#1877)
This PR redesigns 'parameter passing' support for kernel parameters of struct type. All struct type SYCL kernel arguments (except special SYCL types) are now decomposed and their individual fields are passed as separate OpenCL kernel arguments. Following issues are fixed after this PR: 1. Inheritance support for SYCL Kernel 2. Accessors in base class 3. Performance issues due to passing SYCL special types twice. 4. Nested Arrays This PR was a collaborative effort, with patches from Mariya Podchishchaeva and Elizabeth Andrews.
1 parent 768f74f commit 0b2de9e

22 files changed

+757
-185
lines changed

‎clang/lib/Sema/SemaSYCL.cpp‎

Lines changed: 197 additions & 84 deletions
Large diffs are not rendered by default.

‎clang/test/CodeGenSYCL/accessor_inheritance.cpp‎

Lines changed: 94 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,94 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
2+
#include <sycl.hpp>
3+
4+
struct Base {
5+
int A, B;
6+
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> AccField;
7+
};
8+
9+
struct Captured : Base,
10+
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> {
11+
int C;
12+
};
13+
14+
int main() {
15+
Captured Obj;
16+
cl::sycl::kernel_single_task<class kernel>(
17+
[=]() {
18+
Obj.use();
19+
});
20+
return 0;
21+
}
22+
23+
// Check kernel parameters
24+
// CHECK: %[[RANGE_TYPE:"struct.*cl::sycl::range"]]
25+
// CHECK: %[[ID_TYPE:"struct.*cl::sycl::id"]]
26+
// CHECK: define spir_kernel void @_ZTSZ4mainE6kernel
27+
// CHECK-SAME: i32 [[ARG_A:%[a-zA-Z0-9_]+]],
28+
// CHECK-SAME: i32 [[ARG_B:%[a-zA-Z0-9_]+]],
29+
// CHECK-SAME: i8 addrspace(1)* [[ACC1_DATA:%[a-zA-Z0-9_]+]],
30+
// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE1:%[a-zA-Z0-9_]+]],
31+
// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC1_RANGE2:%[a-zA-Z0-9_]+]],
32+
// CHECK-SAME: %[[ID_TYPE]]* byval(%[[ID_TYPE]]) align 4 [[ACC1_ID:%[a-zA-Z0-9_]+]],
33+
// CHECK-SAME: i8 addrspace(1)* [[ACC2_DATA:%[a-zA-Z0-9_]+]],
34+
// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE1:%[a-zA-Z0-9_]+]],
35+
// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC2_RANGE2:%[a-zA-Z0-9_]+]],
36+
// CHECK-SAME: %[[ID_TYPE]]* byval(%[[ID_TYPE]]) align 4 [[ACC2_ID:%[a-zA-Z0-9_]+]],
37+
// CHECK-SAME: i32 [[ARG_C:%[a-zA-Z0-9_]+]])
38+
39+
// Allocas for kernel parameters
40+
// CHECK: [[ARG_A]].addr = alloca i32
41+
// CHECK: [[ARG_B]].addr = alloca i32
42+
// CHECK: [[ACC1_DATA]].addr = alloca i8 addrspace(1)*
43+
// CHECK: [[ACC2_DATA]].addr = alloca i8 addrspace(1)*
44+
// CHECK: [[ARG_C]].addr = alloca i32
45+
//
46+
// Lambda object alloca
47+
// CHECK: [[KERNEL_OBJ:%[a-zA-Z0-9_]+]] = alloca %"class.{{.*}}.anon"
48+
//
49+
// Kernel argument stores
50+
// CHECK: store i32 [[ARG_A]], i32* [[ARG_A]].addr
51+
// CHECK: store i32 [[ARG_B]], i32* [[ARG_B]].addr
52+
// CHECK: store i8 addrspace(1)* [[ACC1_DATA]], i8 addrspace(1)** [[ACC1_DATA]].addr
53+
// CHECK: store i8 addrspace(1)* [[ACC2_DATA]], i8 addrspace(1)** [[ACC2_DATA]].addr
54+
// CHECK: store i32 [[ARG_C]], i32* [[ARG_C]].addr
55+
//
56+
// Check A and B scalar fields initialization
57+
// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0
58+
// CHECK: [[BITCAST:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to %struct{{.*}}Base*
59+
// CHECK: [[FIELD_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 0
60+
// CHECK: [[ARG_A_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32* [[ARG_A]].addr
61+
// CHECK: store i32 [[ARG_A_LOAD]], i32* [[FIELD_A]]
62+
// CHECK: [[FIELD_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 1
63+
// CHECK: [[ARG_B_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32* [[ARG_B]].addr
64+
// CHECK: store i32 [[ARG_B_LOAD]], i32* [[FIELD_B]]
65+
//
66+
// Check accessors initialization
67+
// CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 2
68+
// CHECK: [[ACC1_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
69+
// Default constructor call
70+
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0EEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]])
71+
// CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to i8*
72+
// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 20
73+
// CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8* [[GEP1]] to %"class{{.*}}cl::sycl::accessor"*
74+
// CHECK: [[ACC2_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[BITCAST2]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
75+
// Default constructor call
76+
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0EEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]])
77+
78+
// CHECK C field initialization
79+
// CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured* [[GEP]], i32 0, i32 2
80+
// CHECK: [[ARG_C_LOAD:%[a-zA-Z0-9_]+]] = load i32, i32* [[ARG_C]].addr
81+
// CHECK: store i32 [[ARG_C_LOAD]], i32* [[FIELD_C]]
82+
//
83+
// Check __init method calls
84+
// CHECK: [[GEP2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0
85+
// CHECK: [[BITCAST3:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP2]] to %struct{{.*}}Base*
86+
// CHECK: [[ACC1_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST3]], i32 0, i32 2
87+
// CHECK: [[ACC1_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)** [[ACC1_DATA]].addr
88+
// CHECK: [[ACC1_AS_CAST1:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC1_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
89+
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST1]], i8 addrspace(1)* [[ACC1_DATA_LOAD]]
90+
//
91+
// CHECK: [[GEP3:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon", %"class._ZTSZ4mainE3$_0.anon"* [[KERNEL_OBJ]], i32 0, i32 0
92+
// CHECK: [[ACC2_DATA_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)** [[ACC2_DATA]].addr
93+
// CHECK: [[AS_CAST_CAPTURED:%[a-zA-Z0-9_]+]] = addrspacecast %struct{{.*}}Captured* [[GEP3]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
94+
// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[AS_CAST_CAPTURED]], i8 addrspace(1)* [[ACC2_DATA_LOAD]]

‎clang/test/CodeGenSYCL/inheritance.cpp‎

Lines changed: 84 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,84 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
2+
3+
#include <sycl.hpp>
4+
5+
class second_base {
6+
public:
7+
int e;
8+
};
9+
10+
class InnerFieldBase {
11+
public:
12+
int d;
13+
};
14+
class InnerField : public InnerFieldBase {
15+
int c;
16+
};
17+
18+
struct base {
19+
public:
20+
int b;
21+
InnerField obj;
22+
};
23+
24+
struct derived : base, second_base {
25+
int a;
26+
27+
void operator()() {
28+
}
29+
};
30+
31+
int main() {
32+
cl::sycl::queue q;
33+
34+
q.submit([&](cl::sycl::handler &cgh) {
35+
derived f{};
36+
cgh.single_task(f);
37+
});
38+
39+
return 0;
40+
}
41+
42+
// Check kernel paramters
43+
// CHECK: define spir_kernel void @{{.*}}derived(i32 %_arg_b, i32 %_arg_d, i32 %_arg_c, i32 %_arg_e, i32 %_arg_a)
44+
45+
// Check alloca for kernel paramters
46+
// CHECK: %[[ARG_B:[a-zA-Z0-9_.]+]] = alloca i32, align 4
47+
// CHECK: %[[ARG_D:[a-zA-Z0-9_.]+]] = alloca i32, align 4
48+
// CHECK: %[[ARG_C:[a-zA-Z0-9_.]+]] = alloca i32, align 4
49+
// CHECK: %[[ARG_E:[a-zA-Z0-9_.]+]] = alloca i32, align 4
50+
// CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = alloca i32, align 4
51+
52+
// Check alloca for local functor object
53+
// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = alloca %struct.{{.*}}.derived, align 4
54+
55+
// Initialize field 'b'
56+
// CHECK: %[[BITCAST1:[0-9]+]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to %struct.{{.*}}.base*
57+
// CHECK: %[[GEP_B:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.base, %struct.{{.*}}.base* %[[BITCAST1]], i32 0, i32 0
58+
// CHECK: %[[LOAD_B:[0-9]+]] = load i32, i32* %[[ARG_B]], align 4
59+
// CHECK: store i32 %[[LOAD_B]], i32* %[[GEP_B]], align 4
60+
61+
// Initialize field 'd'
62+
// CHECK: %[[GEP_OBJ:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.base, %struct.{{.*}}.base* %[[BITCAST1]], i32 0, i32 1
63+
// CHECK: %[[BITCAST2:[0-9]+]] = bitcast %class.{{.*}}.InnerField* %[[GEP_OBJ]] to %class.{{.*}}.InnerFieldBase*
64+
// CHECK: %[[GEP_D:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.InnerFieldBase, %class.{{.*}}.InnerFieldBase* %[[BITCAST2]], i32 0, i32 0
65+
// CHECK: %[[LOAD_D:[0-9]+]] = load i32, i32* %[[ARG_D]], align 4
66+
// CHECK: store i32 %[[LOAD_D]], i32* %[[GEP_D]], align 4
67+
68+
// Initialize field 'c'
69+
// CHECK: %[[GEP_C:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.InnerField, %class.{{.*}}.InnerField* %[[GEP_OBJ]], i32 0, i32 1
70+
// CHECK: %[[LOAD_C:[0-9]+]] = load i32, i32* %[[ARG_C]], align 4
71+
// CHECK: store i32 %[[LOAD_C]], i32* %[[GEP_C]], align 4
72+
73+
// Initialize field 'e'
74+
// CHECK: %[[BITCAST3:[0-9]+]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to i8*
75+
// CHECK: %[[GEP_DERIVED:[a-zA-Z0-9]+]] = getelementptr inbounds i8, i8* %[[BITCAST3]], i64 12
76+
// CHECK: %[[BITCAST4:[0-9]+]] = bitcast i8* %[[GEP_DERIVED]] to %class.{{.*}}.second_base*
77+
// CHECK: %[[GEP_E:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.second_base, %class.{{.*}}.second_base* %[[BITCAST4]], i32 0, i32 0
78+
// CHECK: %[[LOAD_E:[0-9]+]] = load i32, i32* %[[ARG_E]], align 4
79+
// CHECK: store i32 %[[LOAD_E]], i32* %[[GEP_E]], align 4
80+
81+
// Initialize field 'a'
82+
// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.derived, %struct.{{.*}}.derived* %[[LOCAL_OBJECT]], i32 0, i32 2
83+
// CHECK: %[[LOAD_A:[0-9]+]] = load i32, i32* %[[ARG_A]], align 4
84+
// CHECK: store i32 %[[LOAD_A]], i32* %[[GEP_A]], align 4

‎clang/test/CodeGenSYCL/integration_header.cpp‎

Lines changed: 20 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
1+
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -emit-llvm
22
// RUN: FileCheck -input-file=%t.h %s
33
//
44
// CHECK: #include <CL/sycl/detail/kernel_desc.hpp>
@@ -28,9 +28,11 @@
2828
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
2929
// CHECK-NEXT: //--- _ZTSZ4mainE12first_kernel
3030
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
31-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 4 },
32-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 16 },
33-
// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 32 },
31+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 },
32+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 8 },
33+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 },
34+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 24 },
35+
// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 40 },
3436
// CHECK-EMPTY:
3537
// CHECK-NEXT: //--- _ZTSN16second_namespace13second_kernelIcEE
3638
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
@@ -46,12 +48,15 @@
4648
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
4749
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 },
4850
// CHECK-EMPTY:
49-
// CHECK-NEXT: //--- _ZTSZ4mainE16accessor_in_base
50-
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 64, 0 },
51-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 8 },
52-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 24 },
53-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 40 },
54-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 52 },
51+
// CHECK-NEXT: //--- _ZTSZ4mainE16accessor_in_base
52+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
53+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 4 },
54+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 8 },
55+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 20 },
56+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 24 },
57+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 36 },
58+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 40 },
59+
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 52 },
5560
// CHECK-EMPTY:
5661
// CHECK-NEXT: };
5762
//
@@ -116,15 +121,13 @@ int main() {
116121
acc2;
117122
int i = 13;
118123
cl::sycl::sampler smplr;
119-
// TODO: Uncomemnt when structures in kernel arguments are correctly processed
120-
// by SYCL compiler
121-
/* struct {
124+
struct {
122125
char c;
123126
int i;
124127
} test_s;
125-
test_s.c = 14;*/
128+
test_s.c = 14;
126129
kernel_single_task<class first_kernel>([=]() {
127-
if (i == 13 /*&& test_s.c == 14*/) {
130+
if (i == 13 && test_s.c == 14) {
128131

129132
acc1.use();
130133
acc2.use();
@@ -151,10 +154,9 @@ int main() {
151154
}
152155
});
153156

154-
// FIXME: We cannot use the member-capture because all the handlers except the
155-
// integration header handler in SemaSYCL don't handle base types right.
156157
accessor_in_base::captured c;
157-
kernel_single_task<class accessor_in_base>([c]() {
158+
kernel_single_task<class accessor_in_base>([=]() {
159+
c.use();
158160
});
159161

160162
return 0;

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

100755100644
File mode changed.

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

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,5 @@
11
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
22
// RUN: FileCheck -input-file=%t.h %s
3-
// XFAIL for now due to : https://github.com/intel/llvm/issues/2018
4-
// XFAIL: *
53

64
// This test checks the integration header when kernel argument
75
// is a struct containing an Accessor array.
@@ -22,7 +20,6 @@
2220
// CHECK: static constexpr
2321
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
2422
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_C
25-
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 24, 0 },
2623
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 },
2724
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 },
2825
// CHECK-EMPTY:

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

Lines changed: 19 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -fsycl-int-header=%t.h -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
2-
// XFAIL: *
32

43
// This test checks a kernel with struct parameter that contains an Accessor array.
54

@@ -29,7 +28,6 @@ int main() {
2928

3029
// CHECK kernel_C parameters
3130
// CHECK: define spir_kernel void @{{.*}}kernel_C
32-
// CHECK-SAME: %struct.{{.*}}.struct_acc_t* byval(%struct.{{.*}}.struct_acc_t) align 4 [[STRUCT:%[a-zA-Z0-9_]+]],
3331
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]],
3432
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
3533
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
@@ -54,32 +52,26 @@ int main() {
5452
// CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
5553
// CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
5654

57-
// Check init of local struct
58-
// CHECK: [[L_STRUCT_ADDR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
59-
// CHECK: [[MEMCPY_DST:%[0-9a-zA-Z_]+]] = bitcast %struct.{{.*}}struct_acc_t* [[L_STRUCT_ADDR]] to i8*
60-
// CHECK: [[MEMCPY_SRC:%[0-9a-zA-Z_]+]] = bitcast %struct.{{.*}}struct_acc_t* %{{[0-9a-zA-Z_]+}} to i8*
61-
// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[MEMCPY_DST]], i8* align 4 [[MEMCPY_SRC]], i64 24, i1 false)
62-
63-
// Check accessor array GEP for member_acc[0]
64-
// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
65-
// CHECK: [[MEMBER1:%[a-zA-Z_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[ACCESSOR_ARRAY1]], i32 0, i32 0
66-
// CHECK: [[Z0:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x %"class.{{.*}}.cl::sycl::accessor"], [2 x %"class.{{.*}}.cl::sycl::accessor"]* [[MEMBER1]], i64 0, i64 0
67-
68-
// Check load from kernel pointer argument alloca
69-
// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}
55+
// Check loop which calls the default constructor for each element of accessor array is emitted.
56+
// CHECK: [[GEP_LAMBDA:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
57+
// CHECK: [[GEP_MEMBER_ACC:%[a-zA-Z_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA]], i32 0, i32 0
58+
// CHECK: [[ARRAY_BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC]], i64 0, i64 0
59+
// CHECK: [[ARRAY_END:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* [[ARRAY_BEGIN]], i64 2
60+
// CHECK: br label %arrayctor.loop
61+
// CHECK: arrayctor.loop:
7062

7163
// Check acc[0] __init method call
72-
// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z0]] to %"class{{.*}}accessor" addrspace(4)*
73-
// 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]])
74-
75-
// Check accessor array GEP for member_acc[1]
76-
// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
77-
// CHECK: [[MEMBER2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[ACCESSOR_ARRAY2]], i32 0, i32 0
78-
// CHECK: [[Z1:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x %"class.{{.*}}.cl::sycl::accessor"], [2 x %"class.{{.*}}.cl::sycl::accessor"]* [[MEMBER2]], i64 0, i64 1
79-
80-
// Check load from kernel pointer argument alloca
81-
// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}
64+
// CHECK: [[GEP_LAMBDA1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
65+
// CHECK: [[GEP_MEMBER_ACC1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA1]], i32 0, i32 0
66+
// CHECK: [[ARRAY_IDX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC1]], i64 0, i64 0
67+
// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr
68+
// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[ARRAY_IDX1]] to [[ACCESSOR]] addrspace(4)*
69+
// CHECK: call spir_func void @{{.*}}__init{{.*}}([[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]])
8270

8371
// Check acc[1] __init method call
84-
// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z1]] to %"class{{.*}}accessor" addrspace(4)*
85-
// 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]])
72+
// CHECK: [[GEP_LAMBDA2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
73+
// CHECK: [[GEP_MEMBER_ACC2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA2]], i32 0, i32 0
74+
// CHECK: [[ARRAY_IDX2:%[a-zA-Z0-9_]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC2]], i64 0, i64 1
75+
// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr
76+
// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[ARRAY_IDX2]] to [[ACCESSOR]] addrspace(4)*
77+
// CHECK: call spir_func void @{{.*}}__init{{.*}}([[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/CodeGenSYCL/kernel-param-pod-array-ih.cpp‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,5 @@
11
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
22
// RUN: FileCheck -input-file=%t.h %s
3-
43
// This test checks the integration header generated for a kernel
54
// with an argument that is a POD array.
65

0 commit comments

Comments
 (0)