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

Commit 1c63fab

Browse files
PennycookPennycook
committed
[SYCL] Enable tests for queue shortcuts
Signed-off-by: John Pennycook <[email protected]>
1 parent 8ca8249 commit 1c63fab

File tree

2 files changed

+125
-59
lines changed

2 files changed

+125
-59
lines changed

‎SYCL/Reduction/reduction_span.cpp

Lines changed: 54 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -39,23 +39,30 @@ size_t getLinearId(range<3> Range, id<3> Id) {
3939
return Id[0] * Range[1] * Range[2] + Id[1] * Range[2] + Id[2];
4040
}
4141

42-
template <size_t N, typename T, typename BinaryOperation, typename Range>
42+
enum class submission_mode {
43+
handler,
44+
queue,
45+
};
46+
47+
template <size_t N, typename T, typename BinaryOperation, typename Range,
48+
submission_mode SubmissionMode>
4349
void test(queue Q, Range Rng, T Identity, T Value) {
4450

4551
// Initialize output to identity value
4652
T *Output = malloc_shared<T>(N, Q);
4753
Q.parallel_for(range<1>{N}, [=](id<1> I) { Output[I] = Identity; }).wait();
4854

4955
// Perform generalized "histogram" with N bins
50-
// TODO: Test Q.parallel_for when code_location is fixed
51-
Q.submit([&](handler &CGH) {
52-
CGH.parallel_for(
53-
Rng, reduction(span<T, N>(Output, N), Identity, BinaryOperation()),
54-
[=](auto It, auto &Reducer) {
55-
size_t Index = getLinearId(Rng, It) % N;
56-
Reducer[Index].combine(Value);
57-
});
58-
}).wait();
56+
auto Redu = reduction(span<T, N>(Output, N), Identity, BinaryOperation());
57+
auto Kern = [=](auto It, auto &Reducer) {
58+
size_t Index = getLinearId(Rng, It) % N;
59+
Reducer[Index].combine(Value);
60+
};
61+
if constexpr (SubmissionMode == submission_mode::handler) {
62+
Q.submit([&](handler &CGH) { CGH.parallel_for(Rng, Redu, Kern); }).wait();
63+
} else /*if (SubmissionMode == submission_mode::queue) */ {
64+
Q.parallel_for(Rng, Redu, Kern).wait();
65+
}
5966

6067
size_t Size = getLinearSize(Rng);
6168

@@ -96,18 +103,43 @@ int main() {
96103

97104
// Tests for small spans that can be privatized efficiently
98105
// Each combination tests a different sycl::reduction implementation
99-
test<16, int, std::plus<int>, sycl::range<1>>(Q, 24, 0, 1);
100-
test<16, float, std::plus<float>, sycl::range<1>>(Q, 24, 0, 1);
101-
test<16, int, std::multiplies<int>, sycl::range<1>>(Q, 24, 1, 2);
102-
test<16, CustomType, CustomBinaryOperation, sycl::range<1>>(
103-
Q, 24, CustomType{0}, CustomType{1});
104-
105-
test<16, int, std::plus<int>, sycl::nd_range<1>>(Q, {24, 8}, 0, 1);
106-
test<16, float, std::plus<float>, sycl::nd_range<1>>(Q, {24, 8}, 0, 1);
107-
test<16, int, std::multiplies<int>, sycl::nd_range<1>>(Q, {24, 8}, 1, 2);
108-
test<16, int, std::bit_or<int>, sycl::nd_range<1>>(Q, {24, 8}, 0, 1);
109-
test<16, CustomType, CustomBinaryOperation, sycl::nd_range<1>>(
110-
Q, {24, 8}, CustomType{0}, CustomType{1});
106+
test<16, int, std::plus<int>, sycl::range<1>, submission_mode::handler>(Q, 24,
107+
0, 1);
108+
test<16, float, std::plus<float>, sycl::range<1>, submission_mode::handler>(
109+
Q, 24, 0, 1);
110+
test<16, int, std::multiplies<int>, sycl::range<1>, submission_mode::handler>(
111+
Q, 24, 1, 2);
112+
test<16, CustomType, CustomBinaryOperation, sycl::range<1>,
113+
submission_mode::handler>(Q, 24, CustomType{0}, CustomType{1});
114+
test<16, int, std::plus<int>, sycl::range<1>, submission_mode::queue>(Q, 24,
115+
0, 1);
116+
test<16, float, std::plus<float>, sycl::range<1>, submission_mode::queue>(
117+
Q, 24, 0, 1);
118+
test<16, int, std::multiplies<int>, sycl::range<1>, submission_mode::queue>(
119+
Q, 24, 1, 2);
120+
test<16, CustomType, CustomBinaryOperation, sycl::range<1>,
121+
submission_mode::queue>(Q, 24, CustomType{0}, CustomType{1});
122+
123+
test<16, int, std::plus<int>, sycl::nd_range<1>, submission_mode::handler>(
124+
Q, {24, 8}, 0, 1);
125+
test<16, float, std::plus<float>, sycl::nd_range<1>,
126+
submission_mode::handler>(Q, {24, 8}, 0, 1);
127+
test<16, int, std::multiplies<int>, sycl::nd_range<1>,
128+
submission_mode::handler>(Q, {24, 8}, 1, 2);
129+
test<16, int, std::bit_or<int>, sycl::nd_range<1>, submission_mode::handler>(
130+
Q, {24, 8}, 0, 1);
131+
test<16, CustomType, CustomBinaryOperation, sycl::nd_range<1>,
132+
submission_mode::handler>(Q, {24, 8}, CustomType{0}, CustomType{1});
133+
test<16, int, std::plus<int>, sycl::nd_range<1>, submission_mode::queue>(
134+
Q, {24, 8}, 0, 1);
135+
test<16, float, std::plus<float>, sycl::nd_range<1>, submission_mode::queue>(
136+
Q, {24, 8}, 0, 1);
137+
test<16, int, std::multiplies<int>, sycl::nd_range<1>,
138+
submission_mode::queue>(Q, {24, 8}, 1, 2);
139+
test<16, int, std::bit_or<int>, sycl::nd_range<1>, submission_mode::queue>(
140+
Q, {24, 8}, 0, 1);
141+
test<16, CustomType, CustomBinaryOperation, sycl::nd_range<1>,
142+
submission_mode::queue>(Q, {24, 8}, CustomType{0}, CustomType{1});
111143

112144
return NumErrors;
113145
}

‎SYCL/Reduction/reduction_span_pack.cpp

Lines changed: 71 additions & 37 deletions
Original file line numberDiff line numberDiff line change
@@ -40,8 +40,14 @@ size_t getLinearId(range<3> Range, id<3> Id) {
4040
return Id[0] * Range[1] * Range[2] + Id[1] * Range[2] + Id[2];
4141
}
4242

43+
enum class submission_mode {
44+
handler,
45+
queue,
46+
};
47+
4348
// Test a span and a regular sum
44-
template <size_t N, typename T, typename BinaryOperation, typename Range>
49+
template <size_t N, typename T, typename BinaryOperation, typename Range,
50+
submission_mode SubmissionMode>
4551
void test1(queue Q, Range Rng, T Identity, T Value) {
4652

4753
// Initialize output to identity value
@@ -51,17 +57,20 @@ void test1(queue Q, Range Rng, T Identity, T Value) {
5157
Q.parallel_for(range<1>{N}, [=](id<1> I) { Output[I] = Identity; }).wait();
5258

5359
// Perform generalized "histogram" with N bins
54-
// TODO: Test Q.parallel_for when code_location is fixed
55-
Q.submit([&](handler &CGH) {
56-
CGH.parallel_for(
57-
Rng, reduction(Sum, plus<>()),
58-
reduction(span<T, N>(Output, N), Identity, BinaryOperation()),
59-
[=](auto It, auto &ScalarReducer, auto &SpanReducer) {
60-
ScalarReducer++;
61-
size_t Index = getLinearId(Rng, It) % N;
62-
SpanReducer[Index].combine(Value);
63-
});
64-
}).wait();
60+
auto ScalarRedu = reduction(Sum, plus<>());
61+
auto SpanRedu = reduction(span<T, N>(Output, N), Identity, BinaryOperation());
62+
auto Kern = [=](auto It, auto &ScalarReducer, auto &SpanReducer) {
63+
ScalarReducer++;
64+
size_t Index = getLinearId(Rng, It) % N;
65+
SpanReducer[Index].combine(Value);
66+
};
67+
if constexpr (SubmissionMode == submission_mode::handler) {
68+
Q.submit([&](handler &CGH) {
69+
CGH.parallel_for(Rng, ScalarRedu, SpanRedu, Kern);
70+
}).wait();
71+
} else /*if (SubmissionMode == submission_mode::queue) */ {
72+
Q.parallel_for(Rng, ScalarRedu, SpanRedu, Kern).wait();
73+
}
6574

6675
size_t Size = getLinearSize(Rng);
6776

@@ -89,7 +98,8 @@ void test1(queue Q, Range Rng, T Identity, T Value) {
8998
}
9099

91100
// Test two spans
92-
template <size_t N, typename T, typename BinaryOperation, typename Range>
101+
template <size_t N, typename T, typename BinaryOperation, typename Range,
102+
submission_mode SubmissionMode>
93103
void test2(queue Q, Range Rng, T Identity, T Value) {
94104

95105
// Initialize output to identity value
@@ -99,17 +109,20 @@ void test2(queue Q, Range Rng, T Identity, T Value) {
99109
Q.parallel_for(range<1>{N}, [=](id<1> I) { Output2[I] = Identity; }).wait();
100110

101111
// Perform generalized "histogram" with N bins
102-
// TODO: Test Q.parallel_for when code_location is fixed
103-
Q.submit([&](handler &CGH) {
104-
CGH.parallel_for(
105-
Rng, reduction(span<int, N>(Output1, N), plus<>()),
106-
reduction(span<T, N>(Output2, N), Identity, BinaryOperation()),
107-
[=](auto It, auto &Reducer1, auto &Reducer2) {
108-
size_t Index = getLinearId(Rng, It) % N;
109-
Reducer1[Index]++;
110-
Reducer2[Index].combine(Value);
111-
});
112-
}).wait();
112+
auto Redu1 = reduction(span<int, N>(Output1, N), plus<>());
113+
auto Redu2 = reduction(span<T, N>(Output2, N), Identity, BinaryOperation());
114+
auto Kern = [=](auto It, auto &Reducer1, auto &Reducer2) {
115+
size_t Index = getLinearId(Rng, It) % N;
116+
Reducer1[Index]++;
117+
Reducer2[Index].combine(Value);
118+
};
119+
if constexpr (SubmissionMode == submission_mode::handler) {
120+
Q.submit([&](handler &CGH) {
121+
CGH.parallel_for(Rng, Redu1, Redu2, Kern);
122+
}).wait();
123+
} else /*if (SubmissionMode == submission_mode::queue) */ {
124+
Q.parallel_for(Rng, Redu1, Redu2, Kern).wait();
125+
}
113126

114127
size_t Size = getLinearSize(Rng);
115128
bool Passed = true;
@@ -165,10 +178,10 @@ struct CustomBinaryOperation {
165178
}
166179
};
167180

168-
template <size_t N, typename T, typename BinaryOperation, typename Range>
181+
template <size_t N, typename T, typename BinaryOperation, typename Range, submission_mode SubmissionMode>
169182
void test(queue Q, Range Rng, T Identity, T Value) {
170-
test1<N, T, BinaryOperation, Range>(Q, Rng, Identity, Value);
171-
test2<N, T, BinaryOperation, Range>(Q, Rng, Identity, Value);
183+
test1<N, T, BinaryOperation, Range, SubmissionMode>(Q, Rng, Identity, Value);
184+
test2<N, T, BinaryOperation, Range, SubmissionMode>(Q, Rng, Identity, Value);
172185
}
173186

174187
int main() {
@@ -177,18 +190,39 @@ int main() {
177190
// Tests for small spans that can be privatized efficiently
178191
// Each combination tests a different sycl::reduction implementation
179192
// TODO: Enable range<> tests once parallel_for accepts pack
180-
/*test<16, int, std::plus<int>, sycl::range<1>>(Q, 24, 0, 1);
181-
test<16, float, std::plus<float>, sycl::range<1>>(Q, 24, 0, 1);
182-
test<16, int, std::multiplies<int>, sycl::range<1>>(Q, 24, 1, 2);
183-
test<16, CustomType, CustomBinaryOperation, sycl::range<1>>(Q, 24,
193+
/*test<16, int, std::plus<int>, sycl::range<1>, submission_mode::handler>(Q,
194+
24, 0, 1); test<16, float, std::plus<float>, sycl::range<1>,
195+
submission_mode::handler>(Q, 24, 0, 1); test<16, int, std::multiplies<int>,
196+
sycl::range<1>, submission_mode::handler>(Q, 24, 1, 2); test<16, CustomType,
197+
CustomBinaryOperation, sycl::range<1>, submission_mode::handler>(Q, 24,
198+
CustomType{0}, CustomType{1});
199+
test<16, int, std::plus<int>, sycl::range<1>, submission_mode::queue>(Q, 24,
200+
0, 1); test<16, float, std::plus<float>, sycl::range<1>,
201+
submission_mode::queue>(Q, 24, 0, 1); test<16, int, std::multiplies<int>,
202+
sycl::range<1>, submission_mode::queue>(Q, 24, 1, 2); test<16, CustomType,
203+
CustomBinaryOperation, sycl::range<1>, submission_mode::queue>(Q, 24,
184204
CustomType{0}, CustomType{1});*/
185205

186-
test<16, int, std::plus<int>, sycl::nd_range<1>>(Q, {24, 8}, 0, 1);
187-
test<16, float, std::plus<float>, sycl::nd_range<1>>(Q, {24, 8}, 0, 1);
188-
test<16, int, std::multiplies<int>, sycl::nd_range<1>>(Q, {24, 8}, 1, 2);
189-
test<16, int, std::bit_or<int>, sycl::nd_range<1>>(Q, {24, 8}, 0, 1);
190-
test<16, CustomType, CustomBinaryOperation, sycl::nd_range<1>>(
191-
Q, {24, 8}, CustomType{0}, CustomType{1});
206+
test<16, int, std::plus<int>, sycl::nd_range<1>, submission_mode::handler>(
207+
Q, {24, 8}, 0, 1);
208+
test<16, float, std::plus<float>, sycl::nd_range<1>,
209+
submission_mode::handler>(Q, {24, 8}, 0, 1);
210+
test<16, int, std::multiplies<int>, sycl::nd_range<1>,
211+
submission_mode::handler>(Q, {24, 8}, 1, 2);
212+
test<16, int, std::bit_or<int>, sycl::nd_range<1>, submission_mode::handler>(
213+
Q, {24, 8}, 0, 1);
214+
test<16, CustomType, CustomBinaryOperation, sycl::nd_range<1>,
215+
submission_mode::handler>(Q, {24, 8}, CustomType{0}, CustomType{1});
216+
test<16, int, std::plus<int>, sycl::nd_range<1>, submission_mode::queue>(
217+
Q, {24, 8}, 0, 1);
218+
test<16, float, std::plus<float>, sycl::nd_range<1>, submission_mode::queue>(
219+
Q, {24, 8}, 0, 1);
220+
test<16, int, std::multiplies<int>, sycl::nd_range<1>,
221+
submission_mode::queue>(Q, {24, 8}, 1, 2);
222+
test<16, int, std::bit_or<int>, sycl::nd_range<1>, submission_mode::queue>(
223+
Q, {24, 8}, 0, 1);
224+
test<16, CustomType, CustomBinaryOperation, sycl::nd_range<1>,
225+
submission_mode::queue>(Q, {24, 8}, CustomType{0}, CustomType{1});
192226

193227
return NumErrors;
194228
}

0 commit comments

Comments
 (0)