Skip to content

Commit d5fb2d9

Browse files
rdeodharrdeodhar
committed
Changed tests to work with current array support.
1 parent af0b0c9 commit d5fb2d9

File tree

5 files changed

+82
-90
lines changed

5 files changed

+82
-90
lines changed

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

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,5 @@
11
// RUN: %clang -I %S/Inputs -fsycl-device-only -Xclang -fsycl-int-header=%t.h %s -c -o %T/kernel.spv
22
// RUN: FileCheck -input-file=%t.h %s
3-
// XFAIL: *
43

54
// This test checks the integration header generated when
65
// the kernel argument is an Accessor array.
@@ -21,7 +20,6 @@
2120
// CHECK: static constexpr
2221
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
2322
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_A
24-
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 24, 0 },
2523
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 },
2624
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 },
2725
// CHECK-EMPTY:

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

Lines changed: 24 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// 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-
// XFAIL: *
32

43
// This test checks a kernel argument that is an Accessor array
54

@@ -26,15 +25,14 @@ int main() {
2625

2726
// Check kernel_A parameters
2827
// CHECK: define spir_kernel void @{{.*}}kernel_A
29-
// CHECK-SAME: %struct.{{.*}}.wrapped_array* byval{{.*}}align 4 %_arg_,
30-
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+_1]],
31-
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_2]],
32-
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_3]],
33-
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_4]],
34-
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG2:%[a-zA-Z0-9_]+_5]],
35-
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+_7]],
36-
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+_8]],
37-
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+_9]])
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]])
3836

3937
// CHECK alloca for pointer arguments
4038
// CHECK: [[MEM_ARG1:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8
@@ -51,30 +49,26 @@ int main() {
5149
// CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
5250
// CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
5351

54-
// Check array initialization
55-
// CHECK: arrayinit.body:
56-
// CHECK: arrayinit.end:
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
5755

58-
// CHECK accessor array GEP for acc[0]
59-
// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
60-
// CHECK: [[Z0:%[a-zA-Z0-9_]*]] = getelementptr inbounds{{.*}}[[ACCESSOR_ARRAY1]], i64 0, i64 0
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]]
6158

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

65-
// CHECK acc[0] __init method call
66-
// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z0]] to %"class{{.*}}accessor" addrspace(4)*
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]])
6763

68-
// 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 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
6967

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

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

77-
// CHECK acc[1] __init method call
78-
// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z1]] to %"class{{.*}}accessor" addrspace(4)*
79-
80-
// 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]])
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]])

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

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,4 @@
11
// RUN: %clang -I %S/Inputs -fsycl-device-only -Xclang -fsycl-int-header=%t.h %s -c -o %T/kernel.spv
2-
// RUN: FileCheck -input-file=%t.h %s
3-
// XFAIL: *
4-
52
// This test checks the integration header generated for a kernel
63
// with an argument that is a POD array.
74

@@ -21,7 +18,11 @@
2118
// CHECK: static constexpr
2219
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
2320
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_B
24-
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 400, 0 },
21+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
22+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 4 },
23+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 8 },
24+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 12 },
25+
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 },
2526
// CHECK-EMPTY:
2627
// CHECK-NEXT: };
2728

@@ -43,7 +44,7 @@ __attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
4344

4445
int main() {
4546

46-
int a[100];
47+
int a[5];
4748

4849
a_kernel<class kernel_B>(
4950
[=]() {

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

Lines changed: 16 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// 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-
// XFAIL: *
32

43
// This test checks a kernel with an argument that is a POD array.
54

@@ -14,26 +13,31 @@ __attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {
1413

1514
int main() {
1615

17-
int a[100];
16+
int a[2];
1817

1918
a_kernel<class kernel_B>(
2019
[=]() {
21-
int local = a[3];
20+
int local = a[1];
2221
});
2322
}
2423

2524
// Check kernel_B parameters
2625
// CHECK: define spir_kernel void @{{.*}}kernel_B
27-
// CHECK-SAME: %struct.{{.*}}.wrapped_array* byval{{.*}}align 4 [[ARG_STRUCT:%[a-zA-Z0-9_]+]]
26+
// CHECK-SAME: i32 [[ELEM_ARG0:%[a-zA-Z0-9_]+]],
27+
// CHECK-SAME: i32 [[ELEM_ARG1:%[a-zA-Z_]+_[0-9]+]])
2828

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

32-
// Check init of local array
33-
// CHECK: [[ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
34-
35-
// CHECK: [[ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.wrapped_array, %struct.{{.*}}.wrapped_array* [[ARG_STRUCT]], i32 0, i32 0
32+
// Check local variables created for parameters
33+
// CHECK: store i32 [[ELEM_ARG0]], i32* [[ELEM_L0:%[a-zA-Z_]+.addr]], align 4
34+
// CHECK: store i32 [[ELEM_ARG1]], i32* [[ELEM_L1:%[a-zA-Z_]+.addr[0-9]*]], align 4
3635

37-
// CHECK: %{{[a-zA-Z0-9._]+}} = getelementptr inbounds [100 x i32], [100 x i32]* [[ARRAY1]], i64 0, i64 0
38-
39-
// CHECK: %{{[a-zA-Z0-9_]+}} = getelementptr inbounds [100 x i32], [100 x i32]* [[ARRAY2]], i64 0, i64
36+
// Check init of local array
37+
// CHECK: [[ARRAY:%[0-9]*]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
38+
// CHECK: [[ARRAY_BEGIN:%[a-zA-Z_.]+]] = getelementptr inbounds [2 x i32], [2 x i32]* [[ARRAY]], i64 0, i64 0
39+
// CHECK: [[ARRAY0:%[0-9]*]] = load i32, i32* [[ELEM_L0]], align 4
40+
// CHECK: store i32 [[ARRAY0]], i32* [[ARRAY_BEGIN]], align 4
41+
// CHECK: [[ARRAY_ELEMENT:%[a-zA-Z_.]+]] = getelementptr inbounds i32, i32* %arrayinit.begin, i64 1
42+
// CHECK: [[ARRAY1:%[0-9]*]] = load i32, i32* [[ELEM_L1]], align 4
43+
// CHECK: store i32 [[ARRAY1]], i32* [[ARRAY_ELEMENT]], align 4

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

Lines changed: 36 additions & 41 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
11
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
2-
// XFAIL: *
32

43
// This test checks that compiler generates correct kernel arguments for
54
// arrays, Accessor arrays, and structs containing Accessors.
@@ -19,7 +18,7 @@ int main() {
1918
accessor<int, 1, access::mode::read_write, access::target::global_buffer>;
2019

2120
Accessor acc[2];
22-
int a[100];
21+
int a[2];
2322
struct struct_acc_t {
2423
Accessor member_acc[4];
2524
} struct_acc;
@@ -31,7 +30,7 @@ int main() {
3130

3231
a_kernel<class kernel_B>(
3332
[=]() {
34-
int local = a[3];
33+
int local = a[1];
3534
});
3635

3736
a_kernel<class kernel_C>(
@@ -41,8 +40,7 @@ int main() {
4140
}
4241

4342
// Check kernel_A parameters
44-
// 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>)'
45-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'wrapped_array'
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>)'
4644
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__global int *'
4745
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'cl::sycl::range<1>'
4846
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'cl::sycl::range<1>'
@@ -56,41 +54,38 @@ int main() {
5654
// CHECK: CXXMemberCallExpr {{.*}} 'void'
5755
// CHECK-NEXT: MemberExpr {{.*}}__init
5856

59-
// CHECK kernel_B parameters
60-
// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (wrapped_array)'
61-
// CHECK-NEXT: ParmVarDecl {{.*}} 'wrapped_array'
62-
// CHECK-NEXT: CompoundStmt
63-
// CHECK-NEXT: DeclStmt
64-
// CHECK-NEXT: VarDecl
65-
// CHECK-NEXT: InitListExpr
66-
// CHECK-NEXT: ArrayInitLoopExpr {{.*}} 'int [100]'
57+
// Check kernel_B parameters
58+
// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (int, int)'
59+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
60+
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
6761

68-
// CHECK kernel_C parameters
69-
// 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>)'
70-
// CHECK-NEXT: ParmVarDecl {{.*}} 'struct {{.*}}'
71-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
72-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
73-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
74-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
75-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
76-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
77-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
78-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
79-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
80-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
81-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
82-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
83-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
84-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
85-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
86-
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
62+
// Correct and enable after struct mebers are extracted into separate parameters
63+
// C HECK kernel_C parameters
64+
// C HECK: 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>)'
65+
// C HECK-NEXT: ParmVarDecl {{.*}} 'struct {{.*}}'
66+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
67+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
68+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
69+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
70+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
71+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
72+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
73+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
74+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
75+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
76+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
77+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
78+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
79+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
80+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
81+
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
8782

88-
// CHECK that four accessor init functions are called
89-
// CHECK: CXXMemberCallExpr {{.*}} 'void'
90-
// CHECK-NEXT: MemberExpr {{.*}}__init
91-
// CHECK: CXXMemberCallExpr {{.*}} 'void'
92-
// CHECK-NEXT: MemberExpr {{.*}}__init
93-
// CHECK: CXXMemberCallExpr {{.*}} 'void'
94-
// CHECK-NEXT: MemberExpr {{.*}}__init
95-
// CHECK: CXXMemberCallExpr {{.*}} 'void'
96-
// CHECK-NEXT: MemberExpr {{.*}}__init
83+
// C HECK that four accessor init functions are called
84+
// C HECK: CXXMemberCallExpr {{.*}} 'void'
85+
// C HECK-NEXT: MemberExpr {{.*}}__init
86+
// C HECK: CXXMemberCallExpr {{.*}} 'void'
87+
// C HECK-NEXT: MemberExpr {{.*}}__init
88+
// C HECK: CXXMemberCallExpr {{.*}} 'void'
89+
// C HECK-NEXT: MemberExpr {{.*}}__init
90+
// C HECK: CXXMemberCallExpr {{.*}} 'void'
91+
// C HECK-NEXT: MemberExpr {{.*}}__init

0 commit comments

Comments
 (0)