Skip to content

Commit b7eabf0

Browse files
srividya-sundaramsrividya-sundaram
authored
[SYCL][NFC] Update FE tests to have a common infrastructure (#2996)
This patch is a first pass at updating FE tests to use the mock "sycl.hpp" header, use SYCL functions for invoking kernels from the mock header (single_task, parallel_for), use sycl:: namespace, add comments, update variable names, update/remove tests etc.
1 parent 084d83a commit b7eabf0

15 files changed

+970
-870
lines changed

‎clang/test/SemaSYCL/Inputs/sycl.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33

44
// Shared code for SYCL tests
55

6-
namespace cl {
6+
inline namespace cl {
77
namespace sycl {
88
namespace access {
99

‎clang/test/SemaSYCL/accessor_inheritance.cpp‎

Lines changed: 35 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -1,65 +1,75 @@
1-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
2-
#include "Inputs/sycl.hpp"
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -sycl-std=2020 %s | FileCheck %s
32

4-
struct Base {
3+
// This test checks inheritance support for struct types with accessors
4+
// passed as kernel arguments, which are decomposed to individual fields.
5+
6+
#include "sycl.hpp"
7+
8+
sycl::queue myQueue;
9+
10+
struct AccessorBase {
511
int A, B;
6-
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> AccField;
12+
sycl::accessor<char, 1, sycl::access::mode::read> AccField;
713
};
814

9-
struct Captured : Base,
10-
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> {
15+
struct AccessorDerived : AccessorBase,
16+
sycl::accessor<char, 1, sycl::access::mode::read> {
1117
int C;
1218
};
1319

1420
int main() {
15-
Captured Obj;
16-
cl::sycl::kernel_single_task<class kernel>(
17-
[=]() {
18-
Obj.use();
19-
});
21+
AccessorDerived DerivedObject;
22+
myQueue.submit([&](sycl::handler &h) {
23+
h.single_task<class kernel>(
24+
[=] {
25+
DerivedObject.use();
26+
});
27+
});
28+
29+
return 0;
2030
}
2131

2232
// Check kernel parameters
23-
// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (int, int, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int)'
33+
// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (int, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int)'
2434
// CHECK: ParmVarDecl{{.*}} used _arg_A 'int'
2535
// CHECK: ParmVarDecl{{.*}} used _arg_B 'int'
2636
// CHECK: ParmVarDecl{{.*}} used _arg_AccField '__global char *'
27-
// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'cl::sycl::range<1>'
28-
// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'cl::sycl::range<1>'
29-
// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'cl::sycl::id<1>'
37+
// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'sycl::range<1>'
38+
// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'sycl::range<1>'
39+
// CHECK: ParmVarDecl{{.*}} used _arg_AccField 'sycl::id<1>'
3040
// CHECK: ParmVarDecl{{.*}} used _arg__base '__global char *'
31-
// CHECK: ParmVarDecl{{.*}} used _arg__base 'cl::sycl::range<1>'
32-
// CHECK: ParmVarDecl{{.*}} used _arg__base 'cl::sycl::range<1>'
33-
// CHECK: ParmVarDecl{{.*}} used _arg__base 'cl::sycl::id<1>'
41+
// CHECK: ParmVarDecl{{.*}} used _arg__base 'sycl::range<1>'
42+
// CHECK: ParmVarDecl{{.*}} used _arg__base 'sycl::range<1>'
43+
// CHECK: ParmVarDecl{{.*}} used _arg__base 'sycl::id<1>'
3444
// CHECK: ParmVarDecl{{.*}} used _arg_C 'int'
3545

3646
// Check lambda initialization
3747
// CHECK: VarDecl {{.*}} used '(lambda at {{.*}}accessor_inheritance.cpp
3848
// CHECK-NEXT: InitListExpr {{.*}}
39-
// CHECK-NEXT: InitListExpr {{.*}} 'Captured'
40-
// CHECK-NEXT: InitListExpr {{.*}} 'Base'
49+
// CHECK-NEXT: InitListExpr {{.*}} 'AccessorDerived'
50+
// CHECK-NEXT: InitListExpr {{.*}} 'AccessorBase'
4151
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
4252
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_A' 'int'
4353
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
4454
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_B' 'int'
45-
// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read>':'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' 'void () noexcept'
46-
// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read>':'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' 'void () noexcept'
55+
// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::accessor<char, 1, sycl::access::mode::read>':'sycl::accessor<char, 1, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' 'void () noexcept'
56+
// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::accessor<char, 1, sycl::access::mode::read>':'sycl::accessor<char, 1, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::false_t>' 'void () noexcept'
4757
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
4858
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_C' 'int'
4959

5060
// Check __init calls
5161
// CHECK: CXXMemberCallExpr {{.*}} 'void'
5262
// CHECK-NEXT: MemberExpr {{.*}} .__init
5363
// CHECK-NEXT: MemberExpr {{.*}} .AccField
54-
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'Base' lvalue <DerivedToBase (Base)>
55-
// CHECK-NEXT: MemberExpr {{.*}} 'Captured' lvalue .
64+
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'AccessorBase' lvalue <DerivedToBase (AccessorBase)>
65+
// CHECK-NEXT: MemberExpr {{.*}} 'AccessorDerived' lvalue .
5666
// CHECK-NEXT: DeclRefExpr {{.*}}'(lambda at {{.*}}accessor_inheritance.cpp
5767
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global char *' <LValueToRValue>
5868
// CHECK-NEXT: DeclRefExpr {{.*}} '__global char *' lvalue ParmVar {{.*}} '_arg_AccField' '__global char *'
5969

6070
// CHECK: CXXMemberCallExpr {{.*}} 'void'
6171
// CHECK-NEXT: MemberExpr{{.*}} lvalue .__init
62-
// CHECK-NEXT: MemberExpr{{.*}}'Captured' lvalue .
72+
// CHECK-NEXT: MemberExpr{{.*}}'AccessorDerived' lvalue .
6373
// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}accessor_inheritance.cpp
6474
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global char *' <LValueToRValue>
6575
// CHECK-NEXT: DeclRefExpr {{.*}} '__global char *' lvalue ParmVar {{.*}} '_arg__base' '__global char *'

‎clang/test/SemaSYCL/accessors-targets-image.cpp‎

Lines changed: 54 additions & 45 deletions
Original file line numberDiff line numberDiff line change
@@ -1,66 +1,75 @@
1-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -sycl-std=2020 %s | FileCheck %s
22

3-
// This test checks that compiler generates correct kernel wrapper arguments for
3+
// This test checks if the compiler generates correct kernel wrapper arguments for
44
// image accessors targets.
55

6-
#include "Inputs/sycl.hpp"
6+
#include "sycl.hpp"
77

8-
using namespace cl::sycl;
9-
10-
template <typename name, typename Func>
11-
__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
12-
kernelFunc();
13-
}
8+
sycl::queue q;
149

1510
int main() {
1611

17-
accessor<int, 1, access::mode::read,
18-
access::target::image, access::placeholder::false_t>
12+
sycl::accessor<int, 1, sycl::access::mode::read,
13+
sycl::access::target::image, sycl::access::placeholder::false_t>
1914
image_acc1d_read;
20-
kernel<class use_image1d_r>(
21-
[=]() {
22-
image_acc1d_read.use();
23-
});
2415

25-
accessor<int, 2, access::mode::read,
26-
access::target::image, access::placeholder::false_t>
16+
q.submit([&](sycl::handler &h) {
17+
h.single_task<class use_image1d_r>(
18+
[=] {
19+
image_acc1d_read.use();
20+
});
21+
});
22+
23+
sycl::accessor<int, 2, sycl::access::mode::read,
24+
sycl::access::target::image, sycl::access::placeholder::false_t>
2725
image_acc2d_read;
28-
kernel<class use_image2d_r>(
29-
[=]() {
30-
image_acc2d_read.use();
31-
});
26+
q.submit([&](sycl::handler &h) {
27+
h.single_task<class use_image2d_r>(
28+
[=] {
29+
image_acc2d_read.use();
30+
});
31+
});
3232

33-
accessor<int, 3, access::mode::read,
34-
access::target::image, access::placeholder::false_t>
33+
sycl::accessor<int, 3, sycl::access::mode::read,
34+
sycl::access::target::image, sycl::access::placeholder::false_t>
3535
image_acc3d_read;
36-
kernel<class use_image3d_r>(
37-
[=]() {
38-
image_acc3d_read.use();
39-
});
4036

41-
accessor<int, 1, access::mode::write,
42-
access::target::image, access::placeholder::false_t>
37+
q.submit([&](sycl::handler &h) {
38+
h.single_task<class use_image3d_r>(
39+
[=] {
40+
image_acc3d_read.use();
41+
});
42+
});
43+
44+
sycl::accessor<int, 1, sycl::access::mode::write,
45+
sycl::access::target::image, sycl::access::placeholder::false_t>
4346
image_acc1d_write;
44-
kernel<class use_image1d_w>(
45-
[=]() {
46-
image_acc1d_write.use();
47-
});
47+
q.submit([&](sycl::handler &h) {
48+
h.single_task<class use_image1d_w>(
49+
[=] {
50+
image_acc1d_write.use();
51+
});
52+
});
4853

49-
accessor<int, 2, access::mode::write,
50-
access::target::image, access::placeholder::false_t>
54+
sycl::accessor<int, 2, sycl::access::mode::write,
55+
sycl::access::target::image, sycl::access::placeholder::false_t>
5156
image_acc2d_write;
52-
kernel<class use_image2d_w>(
53-
[=]() {
54-
image_acc2d_write.use();
55-
});
57+
q.submit([&](sycl::handler &h) {
58+
h.single_task<class use_image2d_w>(
59+
[=] {
60+
image_acc2d_write.use();
61+
});
62+
});
5663

57-
accessor<int, 3, access::mode::write,
58-
access::target::image, access::placeholder::false_t>
64+
sycl::accessor<int, 3, sycl::access::mode::write,
65+
sycl::access::target::image, sycl::access::placeholder::false_t>
5966
image_acc3d_write;
60-
kernel<class use_image3d_w>(
61-
[=]() {
62-
image_acc3d_write.use();
63-
});
67+
q.submit([&](sycl::handler &h) {
68+
h.single_task<class use_image3d_w>(
69+
[=] {
70+
image_acc3d_write.use();
71+
});
72+
});
6473
}
6574

6675
// CHECK: {{.*}}use_image1d_r 'void (__read_only image1d_t)'

‎clang/test/SemaSYCL/accessors-targets.cpp‎

Lines changed: 37 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -1,41 +1,47 @@
1-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -sycl-std=2020 %s | FileCheck %s
22

3-
// This test checks that compiler generates correct kernel wrapper arguments for
3+
// This test checks that the compiler generates correct kernel wrapper arguments for
44
// different accessors targets.
55

6-
#include "Inputs/sycl.hpp"
6+
#include "sycl.hpp"
77

8-
using namespace cl::sycl;
9-
10-
template <typename name, typename Func>
11-
__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
12-
kernelFunc();
13-
}
8+
sycl::queue q;
149

1510
int main() {
16-
17-
accessor<int, 1, access::mode::read_write,
18-
access::target::local>
11+
// Access work-group local memory with read and write access.
12+
sycl::accessor<int, 1, sycl::access::mode::read_write,
13+
sycl::access::target::local>
1914
local_acc;
20-
accessor<int, 1, access::mode::read_write,
21-
access::target::global_buffer>
15+
// Access buffer via global memory with read and write access.
16+
sycl::accessor<int, 1, sycl::access::mode::read_write,
17+
sycl::access::target::global_buffer>
2218
global_acc;
23-
accessor<int, 1, access::mode::read_write,
24-
access::target::constant_buffer>
19+
// Access buffer via constant memory with read and write access.
20+
sycl::accessor<int, 1, sycl::access::mode::read_write,
21+
sycl::access::target::constant_buffer>
2522
constant_acc;
26-
kernel<class use_local>(
27-
[=]() {
28-
local_acc.use();
29-
});
30-
kernel<class use_global>(
31-
[=]() {
32-
global_acc.use();
33-
});
34-
kernel<class use_constant>(
35-
[=]() {
36-
constant_acc.use();
37-
});
23+
24+
q.submit([&](sycl::handler &h) {
25+
h.single_task<class use_local>(
26+
[=] {
27+
local_acc.use();
28+
});
29+
});
30+
31+
q.submit([&](sycl::handler &h) {
32+
h.single_task<class use_global>(
33+
[=] {
34+
global_acc.use();
35+
});
36+
});
37+
38+
q.submit([&](sycl::handler &h) {
39+
h.single_task<class use_constant>(
40+
[=] {
41+
constant_acc.use();
42+
});
43+
});
3844
}
39-
// CHECK: {{.*}}use_local{{.*}} 'void (__local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
40-
// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
41-
// CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
45+
// CHECK: {{.*}}use_local{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
46+
// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
47+
// CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'

‎clang/test/SemaSYCL/allow-constexpr-recursion.cpp‎

Lines changed: 13 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,10 @@
1-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -fcxx-exceptions -Wno-return-type -Wno-sycl-2017-compat -verify -fsyntax-only -std=c++20 -Werror=vla %s
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fcxx-exceptions -Wno-return-type -sycl-std=2020 -verify -fsyntax-only -std=c++20 -Werror=vla %s
22

3-
template <typename name, typename Func>
4-
__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
5-
kernelFunc();
6-
}
3+
// This test verifies that a SYCL kernel executed on a device, cannot call a recursive function.
4+
5+
#include "sycl.hpp"
6+
7+
sycl::queue q;
78

89
// expected-note@+1{{function implemented using recursion declared here}}
910
constexpr int constexpr_recurse1(int n);
@@ -71,6 +72,11 @@ void constexpr_recurse_test_err() {
7172
}
7273

7374
int main() {
74-
kernel_single_task<class fake_kernel>([]() { constexpr_recurse_test(); });
75-
kernel_single_task<class fake_kernel>([]() { constexpr_recurse_test_err(); });
75+
q.submit([&](sycl::handler &h) {
76+
h.single_task<class fake_kernel>([]() { constexpr_recurse_test(); });
77+
});
78+
79+
q.submit([&](sycl::handler &h) {
80+
h.single_task<class fake_kernel>([]() { constexpr_recurse_test_err(); });
81+
});
7682
}

0 commit comments

Comments
 (0)