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

Commit 8df801c

Browse files
maximdimakovmaximdimakov
authored
[SYCL] Add tests for SYCL 2020 function objects (#310)
Signed-off-by: mdimakov <[email protected]>
1 parent 8cb1db6 commit 8df801c

File tree

3 files changed

+453
-0
lines changed

3 files changed

+453
-0
lines changed

‎SYCL/GroupAlgorithm/exclusive_scan_sycl2020.cpp

+180
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,180 @@
1+
// UNSUPPORTED: cuda
2+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %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+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o %t13.out
8+
9+
#include "support.h"
10+
#include <CL/sycl.hpp>
11+
#include <algorithm>
12+
#include <cassert>
13+
#include <limits>
14+
#include <numeric>
15+
#include <vector>
16+
using namespace sycl;
17+
18+
template <class SpecializationKernelName, int TestNumber>
19+
class exclusive_scan_kernel;
20+
21+
// std::exclusive_scan isn't implemented yet, so use serial implementation
22+
// instead
23+
// TODO: use std::exclusive_scan when it will be supported
24+
namespace emu {
25+
template <typename InputIterator, typename OutputIterator,
26+
class BinaryOperation, typename T>
27+
OutputIterator exclusive_scan(InputIterator first, InputIterator last,
28+
OutputIterator result, T init,
29+
BinaryOperation binary_op) {
30+
T partial = init;
31+
for (InputIterator it = first; it != last; ++it) {
32+
*(result++) = partial;
33+
partial = binary_op(partial, *it);
34+
}
35+
return result;
36+
}
37+
} // namespace emu
38+
39+
template <typename SpecializationKernelName, typename InputContainer,
40+
typename OutputContainer, class BinaryOperation>
41+
void test(queue q, InputContainer input, OutputContainer output,
42+
BinaryOperation binary_op,
43+
typename OutputContainer::value_type identity) {
44+
typedef typename InputContainer::value_type InputT;
45+
typedef typename OutputContainer::value_type OutputT;
46+
typedef class exclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
47+
constexpr size_t G = 64;
48+
constexpr size_t N = input.size();
49+
std::vector<OutputT> expected(N);
50+
51+
// checking
52+
// template <typename Group, typename T, typename BinaryOperation>
53+
// T exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op)
54+
{
55+
buffer<InputT> in_buf(input.data(), input.size());
56+
buffer<OutputT> out_buf(output.data(), output.size());
57+
q.submit([&](handler &cgh) {
58+
accessor in{in_buf, cgh, sycl::read_only};
59+
accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
60+
cgh.parallel_for<kernel_name0>(nd_range<1>(G, G), [=](nd_item<1> it) {
61+
group<1> g = it.get_group();
62+
int lid = it.get_local_id(0);
63+
out[lid] = exclusive_scan_over_group(g, in[lid], binary_op);
64+
});
65+
});
66+
}
67+
emu::exclusive_scan(input.begin(), input.begin() + G, expected.begin(),
68+
identity, binary_op);
69+
assert(std::equal(output.begin(), output.begin() + G, expected.begin()));
70+
71+
typedef class exclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
72+
constexpr OutputT init = 42;
73+
74+
// checking
75+
// template <typename Group, typename V, typename T, class BinaryOperation>
76+
// T exclusive_scan_over_group(Group g, V x, T init, BinaryOperation
77+
// binary_op)
78+
{
79+
buffer<InputT> in_buf(input.data(), input.size());
80+
buffer<OutputT> out_buf(output.data(), output.size());
81+
q.submit([&](handler &cgh) {
82+
accessor in{in_buf, cgh, sycl::read_only};
83+
accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
84+
cgh.parallel_for<kernel_name1>(nd_range<1>(G, G), [=](nd_item<1> it) {
85+
group<1> g = it.get_group();
86+
int lid = it.get_local_id(0);
87+
out[lid] = exclusive_scan_over_group(g, in[lid], init, binary_op);
88+
});
89+
});
90+
}
91+
emu::exclusive_scan(input.begin(), input.begin() + G, expected.begin(), init,
92+
binary_op);
93+
assert(std::equal(output.begin(), output.begin() + G, expected.begin()));
94+
95+
typedef class exclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
96+
97+
// checking
98+
// template <typename Group, typename InPtr, typename OutPtr,
99+
// class BinaryOperation>
100+
// OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr
101+
// result,
102+
// BinaryOperation binary_op)
103+
{
104+
buffer<InputT> in_buf(input.data(), input.size());
105+
buffer<OutputT> out_buf(output.data(), output.size());
106+
q.submit([&](handler &cgh) {
107+
accessor in{in_buf, cgh, sycl::read_only};
108+
accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
109+
cgh.parallel_for<kernel_name2>(nd_range<1>(G, G), [=](nd_item<1> it) {
110+
group<1> g = it.get_group();
111+
joint_exclusive_scan(g, in.get_pointer(), in.get_pointer() + N,
112+
out.get_pointer(), binary_op);
113+
});
114+
});
115+
}
116+
emu::exclusive_scan(input.begin(), input.begin() + N, expected.begin(),
117+
identity, binary_op);
118+
assert(std::equal(output.begin(), output.begin() + N, expected.begin()));
119+
120+
typedef class exclusive_scan_kernel<SpecializationKernelName, 3> kernel_name3;
121+
122+
// checking
123+
// template <typename Group, typename InPtr, typename OutPtr, typename T,
124+
// class BinaryOperation>
125+
// OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr
126+
// result, T init,
127+
// BinaryOperation binary_op)
128+
{
129+
buffer<InputT> in_buf(input.data(), input.size());
130+
buffer<OutputT> out_buf(output.data(), output.size());
131+
q.submit([&](handler &cgh) {
132+
accessor in{in_buf, cgh, sycl::read_only};
133+
accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
134+
cgh.parallel_for<kernel_name3>(nd_range<1>(G, G), [=](nd_item<1> it) {
135+
group<1> g = it.get_group();
136+
joint_exclusive_scan(g, in.get_pointer(), in.get_pointer() + N,
137+
out.get_pointer(), init, binary_op);
138+
});
139+
});
140+
}
141+
emu::exclusive_scan(input.begin(), input.begin() + N, expected.begin(), init,
142+
binary_op);
143+
assert(std::equal(output.begin(), output.begin() + N, expected.begin()));
144+
}
145+
146+
int main() {
147+
queue q;
148+
if (!isSupportedDevice(q.get_device())) {
149+
std::cout << "Skipping test\n";
150+
return 0;
151+
}
152+
153+
constexpr int N = 128;
154+
std::array<int, N> input;
155+
std::array<int, N> output;
156+
std::iota(input.begin(), input.end(), 0);
157+
std::fill(output.begin(), output.end(), 0);
158+
159+
test<class KernelNamePlusV>(q, input, output, sycl::plus<>(), 0);
160+
test<class KernelNameMinimumV>(q, input, output, sycl::minimum<>(),
161+
std::numeric_limits<int>::max());
162+
test<class KernelNameMaximumV>(q, input, output, sycl::maximum<>(),
163+
std::numeric_limits<int>::lowest());
164+
165+
test<class KernelNamePlusI>(q, input, output, sycl::plus<int>(), 0);
166+
test<class KernelNameMinimumI>(q, input, output, sycl::minimum<int>(),
167+
std::numeric_limits<int>::max());
168+
test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
169+
std::numeric_limits<int>::lowest());
170+
171+
#ifdef SPIRV_1_3
172+
test<class KernelNameMultipliesI>(q, input, output, sycl::multiplies<int>(),
173+
1);
174+
test<class KernelNameBitOrI>(q, input, output, sycl::bit_or<int>(), 0);
175+
test<class KernelNameBitXorI>(q, input, output, sycl::bit_xor<int>(), 0);
176+
test<class KernelNameBitAndI>(q, input, output, sycl::bit_and<int>(), ~0);
177+
#endif // SPIRV_1_3
178+
179+
std::cout << "Test passed." << std::endl;
180+
}

‎SYCL/GroupAlgorithm/inclusive_scan_sycl2020.cpp

+180
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,180 @@
1+
// UNSUPPORTED: cuda
2+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -I . -o %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+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -DSPIRV_1_3 %s -I . -o %t13.out
8+
9+
#include "support.h"
10+
#include <CL/sycl.hpp>
11+
#include <algorithm>
12+
#include <cassert>
13+
#include <limits>
14+
#include <numeric>
15+
#include <vector>
16+
using namespace sycl;
17+
18+
template <class SpecializationKernelName, int TestNumber>
19+
class inclusive_scan_kernel;
20+
21+
// std::inclusive_scan isn't implemented yet, so use serial implementation
22+
// instead
23+
// TODO: use std::inclusive_scan when it will be supported
24+
namespace emu {
25+
template <typename InputIterator, typename OutputIterator,
26+
class BinaryOperation, typename T>
27+
OutputIterator inclusive_scan(InputIterator first, InputIterator last,
28+
OutputIterator result, BinaryOperation binary_op,
29+
T init) {
30+
T partial = init;
31+
for (InputIterator it = first; it != last; ++it) {
32+
partial = binary_op(partial, *it);
33+
*(result++) = partial;
34+
}
35+
return result;
36+
}
37+
} // namespace emu
38+
39+
template <typename SpecializationKernelName, typename InputContainer,
40+
typename OutputContainer, class BinaryOperation>
41+
void test(queue q, InputContainer input, OutputContainer output,
42+
BinaryOperation binary_op,
43+
typename OutputContainer::value_type identity) {
44+
typedef typename InputContainer::value_type InputT;
45+
typedef typename OutputContainer::value_type OutputT;
46+
typedef class inclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
47+
constexpr size_t N = input.size();
48+
constexpr size_t G = 64;
49+
std::vector<OutputT> expected(N);
50+
51+
// checking
52+
// template <typename Group, typename T, class BinaryOperation>
53+
// T inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op)
54+
{
55+
buffer<InputT> in_buf(input.data(), input.size());
56+
buffer<OutputT> out_buf(output.data(), output.size());
57+
q.submit([&](handler &cgh) {
58+
accessor in{in_buf, cgh, sycl::read_only};
59+
accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
60+
cgh.parallel_for<kernel_name0>(nd_range<1>(G, G), [=](nd_item<1> it) {
61+
group<1> g = it.get_group();
62+
int lid = it.get_local_id(0);
63+
out[lid] = inclusive_scan_over_group(g, in[lid], binary_op);
64+
});
65+
});
66+
}
67+
emu::inclusive_scan(input.begin(), input.begin() + G, expected.begin(),
68+
binary_op, identity);
69+
assert(std::equal(output.begin(), output.begin() + G, expected.begin()));
70+
71+
typedef class inclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
72+
constexpr OutputT init = 42;
73+
74+
// checking
75+
// template <typename Group, typename V, class BinaryOperation, typename T>
76+
// T inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T
77+
// init)
78+
{
79+
buffer<InputT> in_buf(input.data(), input.size());
80+
buffer<OutputT> out_buf(output.data(), output.size());
81+
q.submit([&](handler &cgh) {
82+
accessor in{in_buf, cgh, sycl::read_only};
83+
accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
84+
cgh.parallel_for<kernel_name1>(nd_range<1>(G, G), [=](nd_item<1> it) {
85+
group<1> g = it.get_group();
86+
int lid = it.get_local_id(0);
87+
out[lid] = inclusive_scan_over_group(g, in[lid], binary_op, init);
88+
});
89+
});
90+
}
91+
emu::inclusive_scan(input.begin(), input.begin() + G, expected.begin(),
92+
binary_op, init);
93+
assert(std::equal(output.begin(), output.begin() + G, expected.begin()));
94+
95+
typedef class inclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
96+
97+
// checking
98+
// template <typename Group, typename InPtr, typename OutPtr,
99+
// class BinaryOperation>
100+
// OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr
101+
// result,
102+
// BinaryOperation binary_op)
103+
{
104+
buffer<InputT> in_buf(input.data(), input.size());
105+
buffer<OutputT> out_buf(output.data(), output.size());
106+
q.submit([&](handler &cgh) {
107+
accessor in{in_buf, cgh, sycl::read_only};
108+
accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
109+
cgh.parallel_for<kernel_name2>(nd_range<1>(G, G), [=](nd_item<1> it) {
110+
group<1> g = it.get_group();
111+
joint_inclusive_scan(g, in.get_pointer(), in.get_pointer() + N,
112+
out.get_pointer(), binary_op);
113+
});
114+
});
115+
}
116+
emu::inclusive_scan(input.begin(), input.begin() + N, expected.begin(),
117+
binary_op, identity);
118+
assert(std::equal(output.begin(), output.begin() + N, expected.begin()));
119+
120+
typedef class inclusive_scan_kernel<SpecializationKernelName, 3> kernel_name3;
121+
122+
// checking
123+
// template <typename Group, typename InPtr, typename OutPtr,
124+
// class BinaryOperation, typename T>
125+
// OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr
126+
// result,
127+
// BinaryOperation binary_op, T init)
128+
{
129+
buffer<InputT> in_buf(input.data(), input.size());
130+
buffer<OutputT> out_buf(output.data(), output.size());
131+
q.submit([&](handler &cgh) {
132+
accessor in{in_buf, cgh, sycl::read_only};
133+
accessor out{out_buf, cgh, sycl::write_only, sycl::no_init};
134+
cgh.parallel_for<kernel_name3>(nd_range<1>(G, G), [=](nd_item<1> it) {
135+
group<1> g = it.get_group();
136+
joint_inclusive_scan(g, in.get_pointer(), in.get_pointer() + N,
137+
out.get_pointer(), binary_op, init);
138+
});
139+
});
140+
}
141+
emu::inclusive_scan(input.begin(), input.begin() + N, expected.begin(),
142+
binary_op, init);
143+
assert(std::equal(output.begin(), output.begin() + N, expected.begin()));
144+
}
145+
146+
int main() {
147+
queue q;
148+
if (!isSupportedDevice(q.get_device())) {
149+
std::cout << "Skipping test\n";
150+
return 0;
151+
}
152+
153+
constexpr int N = 128;
154+
std::array<int, N> input;
155+
std::array<int, N> output;
156+
std::iota(input.begin(), input.end(), 0);
157+
std::fill(output.begin(), output.end(), 0);
158+
159+
test<class KernelNamePlusV>(q, input, output, sycl::plus<>(), 0);
160+
test<class KernelNameMinimumV>(q, input, output, sycl::minimum<>(),
161+
std::numeric_limits<int>::max());
162+
test<class KernelNameMaximumV>(q, input, output, sycl::maximum<>(),
163+
std::numeric_limits<int>::lowest());
164+
165+
test<class KernelNamePlusI>(q, input, output, sycl::plus<int>(), 0);
166+
test<class KernelNameMinimumI>(q, input, output, sycl::minimum<int>(),
167+
std::numeric_limits<int>::max());
168+
test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
169+
std::numeric_limits<int>::lowest());
170+
171+
#ifdef SPIRV_1_3
172+
test<class KernelNameMultipliesI>(q, input, output, sycl::multiplies<int>(),
173+
1);
174+
test<class KernelNameBitOrI>(q, input, output, sycl::bit_or<int>(), 0);
175+
test<class KernelNameBitXorI>(q, input, output, sycl::bit_xor<int>(), 0);
176+
test<class KernelNameBitAndI>(q, input, output, sycl::bit_and<int>(), ~0);
177+
#endif // SPIRV_1_3
178+
179+
std::cout << "Test passed." << std::endl;
180+
}

0 commit comments

Comments
 (0)