Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit 4213591

Browse files
vladimirlazvladimirlaz
committed
[SYCL] add functor tests
1 parent 3b3e023 commit 4213591

File tree

2 files changed

+239
-0
lines changed

2 files changed

+239
-0
lines changed

‎SYCL/Basic/Functor/functor_inheritance.cpp

Lines changed: 60 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,60 @@
1+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -o %t.out %s
2+
// RUN: %HOST_RUN_PLACEHOLDER %t.out
3+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
4+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
5+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
6+
7+
#include <CL/sycl.hpp>
8+
#include <cassert>
9+
10+
constexpr auto sycl_read_write = cl::sycl::access::mode::read_write;
11+
constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer;
12+
13+
struct SecondBase {
14+
SecondBase(int _E) : E(_E) {}
15+
int E;
16+
};
17+
18+
struct InnerFieldBase {
19+
InnerFieldBase(int _D) : D(_D) {}
20+
int D;
21+
};
22+
23+
struct InnerField : public InnerFieldBase {
24+
InnerField(int _C, int _D) : C(_C), InnerFieldBase(_D) {}
25+
int C;
26+
};
27+
28+
struct Base {
29+
Base(int _B, int _C, int _D) : B(_B), InnerObj(_C, _D) {}
30+
int B;
31+
InnerField InnerObj;
32+
};
33+
34+
struct Derived : public Base, public SecondBase {
35+
Derived(int _A, int _B, int _C, int _D, int _E,
36+
cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> &_Acc)
37+
: A(_A), Acc(_Acc), /*Out(_Out),*/ Base(_B, _C, _D), SecondBase(_E) {}
38+
void operator()() const {
39+
Acc[0] = this->A + this->B + this->InnerObj.C + this->InnerObj.D + this->E;
40+
}
41+
42+
int A;
43+
cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> Acc;
44+
};
45+
46+
int main() {
47+
int A[] = {10};
48+
{
49+
cl::sycl::queue Q;
50+
cl::sycl::buffer<int, 1> Buf(A, 1);
51+
52+
Q.submit([&](cl::sycl::handler &cgh) {
53+
auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
54+
Derived F = {1, 2, 3, 4, 5, Acc /*, Out*/};
55+
cgh.single_task(F);
56+
});
57+
}
58+
assert(A[0] == 15);
59+
return 0;
60+
}

‎SYCL/Basic/Functor/kernel_functor.cpp

Lines changed: 179 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,179 @@
1+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -o %t.out %s
2+
// RUN: cd %T
3+
// RUN: %HOST_RUN_PLACEHOLDER %t.out
4+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
5+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
6+
7+
//==--- kernel_functor.cpp - Functors as SYCL kernel test ------------------==//
8+
//
9+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
10+
// See https://llvm.org/LICENSE.txt for license information.
11+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
12+
//
13+
//===----------------------------------------------------------------------===//
14+
15+
#include <CL/sycl.hpp>
16+
17+
#include <cassert>
18+
19+
constexpr auto sycl_read_write = cl::sycl::access::mode::read_write;
20+
constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer;
21+
22+
// 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+
// - functor class is defined in a namespace
45+
// - the '()' operator:
46+
// * does not have parameters (to be used in 'single_task').
47+
// * has the 'const' qualifier
48+
namespace ns {
49+
class Functor2 {
50+
public:
51+
Functor2(
52+
int X_,
53+
cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> &Acc_)
54+
: X(X_), Acc(Acc_) {}
55+
56+
// cl::sycl::accessor's operator [] is const, hence 'const' is possible below
57+
void operator()() const { Acc[0] += X; }
58+
59+
private:
60+
int X;
61+
cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> Acc;
62+
};
63+
} // namespace ns
64+
65+
// Case 3:
66+
// - functor class is templated and defined in the translation unit scope
67+
// - the '()' operator:
68+
// * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for').
69+
// * has the 'const' qualifier
70+
template <typename T> class TmplFunctor {
71+
public:
72+
TmplFunctor(
73+
T X_, cl::sycl::accessor<T, 1, sycl_read_write, sycl_global_buffer> &Acc_)
74+
: X(X_), Acc(Acc_) {}
75+
76+
void operator()(cl::sycl::id<1> id) const { Acc[id] += X; }
77+
78+
private:
79+
T X;
80+
cl::sycl::accessor<T, 1, sycl_read_write, sycl_global_buffer> Acc;
81+
};
82+
83+
// Case 4:
84+
// - functor class is templated and defined in the translation unit scope
85+
// - the '()' operator:
86+
// * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for').
87+
// * has the 'const' qualifier
88+
template <typename T> class TmplConstFunctor {
89+
public:
90+
TmplConstFunctor(
91+
T X_, cl::sycl::accessor<T, 1, sycl_read_write, sycl_global_buffer> &Acc_)
92+
: X(X_), Acc(Acc_) {}
93+
94+
void operator()(cl::sycl::id<1> id) const { Acc[id] += X; }
95+
96+
private:
97+
T X;
98+
cl::sycl::accessor<T, 1, sycl_read_write, sycl_global_buffer> Acc;
99+
};
100+
101+
// Exercise non-templated functors in 'single_task'.
102+
int foo(int X) {
103+
int A[] = {10};
104+
{
105+
cl::sycl::queue Q;
106+
cl::sycl::buffer<int, 1> Buf(A, 1);
107+
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+
Q.submit([&](cl::sycl::handler &cgh) {
115+
auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
116+
ns::Functor2 F(X, Acc);
117+
118+
cgh.single_task(F);
119+
});
120+
Q.submit([&](cl::sycl::handler &cgh) {
121+
auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
122+
ns::Functor2 F(X, Acc);
123+
124+
cgh.single_task(F);
125+
});
126+
}
127+
return A[0];
128+
}
129+
130+
#define ARR_LEN(x) sizeof(x) / sizeof(x[0])
131+
132+
// Exercise templated functors in 'parallel_for'.
133+
template <typename T> T bar(T X) {
134+
T A[] = {(T)10, (T)10};
135+
{
136+
cl::sycl::queue Q;
137+
cl::sycl::buffer<T, 1> Buf(A, ARR_LEN(A));
138+
139+
Q.submit([&](cl::sycl::handler &cgh) {
140+
auto Acc =
141+
Buf.template get_access<sycl_read_write, sycl_global_buffer>(cgh);
142+
TmplFunctor<T> F(X, Acc);
143+
144+
cgh.parallel_for(cl::sycl::range<1>(ARR_LEN(A)), F);
145+
});
146+
// Spice with lambdas to make sure functors and lambdas work together.
147+
Q.submit([&](cl::sycl::handler &cgh) {
148+
auto Acc =
149+
Buf.template get_access<sycl_read_write, sycl_global_buffer>(cgh);
150+
cgh.parallel_for<class LambdaKernel>(
151+
cl::sycl::range<1>(ARR_LEN(A)),
152+
[=](cl::sycl::id<1> id) { Acc[id] += X; });
153+
});
154+
Q.submit([&](cl::sycl::handler &cgh) {
155+
auto Acc =
156+
Buf.template get_access<sycl_read_write, sycl_global_buffer>(cgh);
157+
TmplConstFunctor<T> F(X, Acc);
158+
159+
cgh.parallel_for(cl::sycl::range<1>(ARR_LEN(A)), F);
160+
});
161+
}
162+
T res = (T)0;
163+
164+
for (int i = 0; i < ARR_LEN(A); i++)
165+
res += A[i];
166+
return res;
167+
}
168+
169+
int main() {
170+
const int Res1 = foo(10);
171+
const int Res2 = bar(10);
172+
const int Gold1 = 40;
173+
const int Gold2 = 80;
174+
175+
assert(Res1 == Gold1);
176+
assert(Res2 == Gold2);
177+
178+
return 0;
179+
}

0 commit comments

Comments
 (0)