Skip to content

Support for arrays as kernel parameters. #1841

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 21 commits into from
Jul 2, 2020
Merged

Support for arrays as kernel parameters. #1841

Show file tree
Hide file tree
Changes from 17 commits
Commits
Show all changes
21 commits
Select commit Hold shift + click to select a range
7974c35
Support for arrays as kernel parameters.
rdeodhar Jun 9, 2020
4907194
Reusing some memberexpr building code.
rdeodhar Jun 9, 2020
d54c0ca
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 9, 2020
546c58d
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 10, 2020
52ce3f2
Updated support for arrays.
rdeodhar Jun 12, 2020
983b3d5
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 12, 2020
1bf0903
Formatting changes.
rdeodhar Jun 12, 2020
5d5121b
Formatting changes.
rdeodhar Jun 12, 2020
f03edd9
Correction to a test.
rdeodhar Jun 15, 2020
d87b2cc
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 22, 2020
0412db3
Array elements are now passed as individual parameters.
rdeodhar Jun 25, 2020
810af7b
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 25, 2020
00c082f
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 25, 2020
af0b0c9
Corrections to temporarily disable tests expected to fail.
rdeodhar Jun 25, 2020
d5fb2d9
Changed tests to work with current array support.
rdeodhar Jun 26, 2020
db492bd
Decomposed array elements, and changed manner of array element initia…
rdeodhar Jun 27, 2020
59cabac
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 27, 2020
4afc3a3
Removed one redundant check.
rdeodhar Jun 29, 2020
9196a30
Changed how some lit tests are run.
rdeodhar Jun 30, 2020
5660269
Update clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp
rdeodhar Jun 30, 2020
81ace26
Fixed formatting.
rdeodhar Jun 30, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
243 changes: 133 additions & 110 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop

Large diffs are not rendered by default.

7 changes: 2 additions & 5 deletions clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@

// CHECK: static constexpr
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_A
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_A
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 },
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 },
// CHECK-EMPTY:
Expand Down Expand Up @@ -48,8 +48,5 @@ int main() {

Accessor acc[2];

a_kernel<class kernel_A>(
[=]() {
acc[1].use();
});
a_kernel<class kernel_A>([=]() { acc[1].use(); });
}
42 changes: 23 additions & 19 deletions clang/test/CodeGenSYCL/kernel-param-acc-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,41 +34,45 @@ int main() {
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+_7]],
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+_8]])

// Check alloca for pointer arguments
// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8
// CHECK: [[MEM_ARG1]].addr{{[0-9]*}} = alloca i32 addrspace(1)*, align 8
// CHECK alloca for pointer arguments
// CHECK: [[MEM_ARG1:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8
// CHECK: [[MEM_ARG2:%[a-zA-Z0-9_.]+]] = alloca i32 addrspace(1)*, align 8

// Check lambda object alloca
// CHECK lambda object alloca
// CHECK: [[LOCAL_OBJECT:%0]] = alloca %"class.{{.*}}.anon", align 4

// Check allocas for ranges
// CHECK allocas for ranges
// CHECK: [[ACC_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
// CHECK: [[MEM_RANGE1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
// CHECK: [[OFFSET1:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"
// CHECK: [[ACC_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
// CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range"
// CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id"

// Check accessor array GEP for acc[0]
// CHECK accessor array default inits
// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
// CHECK: [[Z0:%[a-zA-Z0-9_]*]] = getelementptr inbounds{{.*}}[[ACCESSOR_ARRAY1]], i64 0, i64 0
// CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY1]], i64 0, i64 0
// CHECK: [[END:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR:.*]], [[ACCESSOR]]* [[BEGIN]], i64 2
// CHECK: [[NEXT0:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1
// CHECK: [[ELEMENT:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1
// CHECK: [[ELEMENT:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 2
// CHECK: [[NEXT1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1

// Check load from kernel pointer argument alloca
// CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}
// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
// CHECK: [[INDEX:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY2]], i64 0, i64 0

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

// 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]])
// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* {{.*}} to [[ACCESSOR]] addrspace(4)*

// Check accessor array GEP for acc[1]
// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
// CHECK: [[Z1:%[a-zA-Z0-9_]*]] = getelementptr inbounds{{.*}}[[ACCESSOR_ARRAY2]], i64 0, i64 1
// CHECK acc[0] __init method call
// 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]])

// Check load from kernel pointer argument alloca
// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]].addr{{[0-9]*}}
// CHECK load from kernel pointer argument alloca
// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG2]]

// Check acc[1] __init method call
// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast %"class{{.*}}accessor"* [[Z1]] to %"class{{.*}}accessor" addrspace(4)*
// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* {{.*}} to [[ACCESSOR]] addrspace(4)*

// CHECK acc[1] __init method call
// 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]])
1 change: 1 addition & 0 deletions clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,6 @@
// RUN: %clang -I %S/Inputs -fsycl-device-only -Xclang -fsycl-int-header=%t.h %s -c -o %T/kernel.spv
// RUN: FileCheck -input-file=%t.h %s
// XFAIL: *

// This test checks the integration header when kernel argument
// is a struct containing an Accessor array.
Expand Down
1 change: 1 addition & 0 deletions clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,5 @@
// 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
// XFAIL: *

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

Expand Down
10 changes: 6 additions & 4 deletions clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,4 @@
// RUN: %clang -I %S/Inputs -fsycl-device-only -Xclang -fsycl-int-header=%t.h %s -c -o %T/kernel.spv
// RUN: FileCheck -input-file=%t.h %s

// This test checks the integration header generated for a kernel
// with an argument that is a POD array.

Expand All @@ -20,7 +18,11 @@
// CHECK: static constexpr
// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
// CHECK-NEXT: //--- _ZTSZ4mainE8kernel_B
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 400, 0 },
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 },
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 4 },
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 8 },
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 12 },
// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 },
// CHECK-EMPTY:
// CHECK-NEXT: };

Expand All @@ -42,7 +44,7 @@ __attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {

int main() {

int a[100];
int a[5];

a_kernel<class kernel_B>(
[=]() {
Expand Down
27 changes: 16 additions & 11 deletions clang/test/CodeGenSYCL/kernel-param-pod-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,26 +13,31 @@ __attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) {

int main() {

int a[100];
int a[2];

a_kernel<class kernel_B>(
[=]() {
int local = a[3];
int local = a[1];
});
}

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

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

// Check init of local array
// CHECK: [[ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0

// CHECK: [[ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.wrapped_array, %struct.{{.*}}.wrapped_array* [[ARG_STRUCT]], i32 0, i32 0
// Check local variables created for parameters
// CHECK: store i32 [[ELEM_ARG0]], i32* [[ELEM_L0:%[a-zA-Z_]+.addr]], align 4
// CHECK: store i32 [[ELEM_ARG1]], i32* [[ELEM_L1:%[a-zA-Z_]+.addr[0-9]*]], align 4

// CHECK: %{{[a-zA-Z0-9._]+}} = getelementptr inbounds [100 x i32], [100 x i32]* [[ARRAY1]], i64 0, i64 0

// CHECK: %{{[a-zA-Z0-9_]+}} = getelementptr inbounds [100 x i32], [100 x i32]* [[ARRAY2]], i64 0, i64
// Check init of local array
// CHECK: [[ARRAY:%[0-9]*]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0
// CHECK: [[ARRAY_BEGIN:%[a-zA-Z_.]+]] = getelementptr inbounds [2 x i32], [2 x i32]* [[ARRAY]], i64 0, i64 0
// CHECK: [[ARRAY0:%[0-9]*]] = load i32, i32* [[ELEM_L0]], align 4
// CHECK: store i32 [[ARRAY0]], i32* [[ARRAY_BEGIN]], align 4
// CHECK: [[ARRAY_ELEMENT:%[a-zA-Z_.]+]] = getelementptr inbounds i32, i32* %arrayinit.begin, i64 1
// CHECK: [[ARRAY1:%[0-9]*]] = load i32, i32* [[ELEM_L1]], align 4
// CHECK: store i32 [[ARRAY1]], i32* [[ARRAY_ELEMENT]], align 4
13 changes: 0 additions & 13 deletions clang/test/SemaSYCL/array-kernel-param-neg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,21 +4,12 @@
// an array of non-trivially copyable structs as SYCL kernel parameter or
// a non-constant size array.

struct A {
int i;
};

struct B {
int i;
B(int _i) : i(_i) {}
B(const B &x) : i(x.i) {}
};

struct C : A {
const A C2;
C() : A{0}, C2{2} {}
};

struct D {
int i;
~D();
Expand All @@ -38,16 +29,12 @@ __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
}

void test() {
A cs1[10];
B nsl1[4] = {1, 2, 3, 4};
C cs2[6];
D nsl2[5];
E es;
kernel_single_task<class kernel_capture_refs>([=] {
int a = cs1[6].i;
// expected-error@+1 {{kernel parameter has non-trivially copy constructible class/struct type}}
int b = nsl1[2].i;
int c = cs2[0].i;
// expected-error@+1 {{kernel parameter has non-trivially destructible class/struct type}}
int d = nsl2[4].i;
});
Expand Down
75 changes: 41 additions & 34 deletions clang/test/SemaSYCL/array-kernel-param.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@ int main() {
accessor<int, 1, access::mode::read_write, access::target::global_buffer>;

Accessor acc[2];
int a[100];
int a[2];
struct struct_acc_t {
Accessor member_acc[4];
} struct_acc;
Expand All @@ -30,7 +30,7 @@ int main() {

a_kernel<class kernel_B>(
[=]() {
int local = a[3];
int local = a[1];
});

a_kernel<class kernel_C>(
Expand All @@ -55,40 +55,47 @@ int main() {
// CHECK-NEXT: MemberExpr {{.*}}__init

// Check kernel_B parameters
// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (wrapped_array)'
// CHECK-NEXT: ParmVarDecl {{.*}} 'wrapped_array'
// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (int, int)'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
// Check kernel_B inits
// CHECK-NEXT: CompoundStmt
// CHECK-NEXT: DeclStmt
// CHECK-NEXT: VarDecl
// CHECK-NEXT: VarDecl {{.*}} cinit
// CHECK-NEXT: InitListExpr
// CHECK-NEXT: ArrayInitLoopExpr {{.*}} 'int [100]'
// CHECK-NEXT: InitListExpr {{.*}} 'int [2]'
// CHECK: ImplicitCastExpr
// CHECK: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
// CHECK: ImplicitCastExpr
// CHECK: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'

// Check kernel_C parameters
// 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>)'
// CHECK-NEXT: ParmVarDecl {{.*}} 'struct {{.*}}'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
// Correct and enable after struct members are extracted into separate parameters
// C HECK kernel_C parameters
// 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>)'
// C HECK-NEXT: ParmVarDecl {{.*}} 'struct {{.*}}'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::range<1>'
// C HECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'cl::sycl::id<1>'

// Check that four accessor init functions are called
// CHECK: CXXMemberCallExpr {{.*}} 'void'
// CHECK-NEXT: MemberExpr {{.*}}__init
// CHECK: CXXMemberCallExpr {{.*}} 'void'
// CHECK-NEXT: MemberExpr {{.*}}__init
// CHECK: CXXMemberCallExpr {{.*}} 'void'
// CHECK-NEXT: MemberExpr {{.*}}__init
// CHECK: CXXMemberCallExpr {{.*}} 'void'
// CHECK-NEXT: MemberExpr {{.*}}__init
// C HECK that four accessor init functions are called
// C HECK: CXXMemberCallExpr {{.*}} 'void'
// C HECK-NEXT: MemberExpr {{.*}}__init
// C HECK: CXXMemberCallExpr {{.*}} 'void'
// C HECK-NEXT: MemberExpr {{.*}}__init
// C HECK: CXXMemberCallExpr {{.*}} 'void'
// C HECK-NEXT: MemberExpr {{.*}}__init
// C HECK: CXXMemberCallExpr {{.*}} 'void'
// C HECK-NEXT: MemberExpr {{.*}}__init
3 changes: 2 additions & 1 deletion sycl/doc/CompilerAndRuntimeDesign.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -168,7 +168,8 @@ __kernel KernelName(global int* a) {
```

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

### SYCL support in the driver

Expand Down
Loading