4
4
// RUN: %CPU_RUN_PLACEHOLDER %t.out
5
5
// RUN: %GPU_RUN_PLACEHOLDER %t.out
6
6
7
- // ==--- kernel_functor.cpp - Functors as SYCL kernel test ------------------==//
7
+ // ==--- kernel_functor.cpp -
8
+ // This test illustrates defining kernels as named function objects (functors)
8
9
//
9
10
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
10
11
// See https://llvm.org/LICENSE.txt for license information.
@@ -20,27 +21,6 @@ constexpr auto sycl_read_write = cl::sycl::access::mode::read_write;
20
21
constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer;
21
22
22
23
// Case 1:
23
- // - functor class is defined in an anonymous namespace
24
- // - the '()' operator:
25
- // * does not have parameters (to be used in 'single_task').
26
- // * has the 'const' qualifier
27
- namespace {
28
- class Functor1 {
29
- public:
30
- Functor1 (
31
- int X_,
32
- cl::sycl::accessor<int , 1 , sycl_read_write, sycl_global_buffer> &Acc_)
33
- : X(X_), Acc(Acc_) {}
34
-
35
- void operator ()() const { Acc[0 ] += X; }
36
-
37
- private:
38
- int X;
39
- cl::sycl::accessor<int , 1 , sycl_read_write, sycl_global_buffer> Acc;
40
- };
41
- } // namespace
42
-
43
- // Case 2:
44
24
// - functor class is defined in a namespace
45
25
// - the '()' operator:
46
26
// * does not have parameters (to be used in 'single_task').
@@ -62,7 +42,7 @@ class Functor2 {
62
42
};
63
43
} // namespace ns
64
44
65
- // Case 3 :
45
+ // Case 2 :
66
46
// - functor class is templated and defined in the translation unit scope
67
47
// - the '()' operator:
68
48
// * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for').
@@ -80,7 +60,7 @@ template <typename T> class TmplFunctor {
80
60
cl::sycl::accessor<T, 1 , sycl_read_write, sycl_global_buffer> Acc;
81
61
};
82
62
83
- // Case 4 :
63
+ // Case 3 :
84
64
// - functor class is templated and defined in the translation unit scope
85
65
// - the '()' operator:
86
66
// * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for').
@@ -105,12 +85,6 @@ int foo(int X) {
105
85
cl::sycl::queue Q;
106
86
cl::sycl::buffer<int , 1 > Buf (A, 1 );
107
87
108
- Q.submit ([&](cl::sycl::handler &cgh) {
109
- auto Acc = Buf.get_access <sycl_read_write, sycl_global_buffer>(cgh);
110
- Functor1 F (X, Acc);
111
-
112
- cgh.single_task (F);
113
- });
114
88
Q.submit ([&](cl::sycl::handler &cgh) {
115
89
auto Acc = Buf.get_access <sycl_read_write, sycl_global_buffer>(cgh);
116
90
ns::Functor2 F (X, Acc);
@@ -169,7 +143,7 @@ template <typename T> T bar(T X) {
169
143
int main () {
170
144
const int Res1 = foo (10 );
171
145
const int Res2 = bar (10 );
172
- const int Gold1 = 40 ;
146
+ const int Gold1 = 30 ;
173
147
const int Gold2 = 80 ;
174
148
175
149
assert (Res1 == Gold1);
0 commit comments