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

Commit 89ae917

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Add identityless reduction test cases (#1640)
This commit adds test cases to existing tests for testing reductions where the identity is unspecified and unknown. --------- Signed-off-by: Larsen, Steffen <[email protected]>
1 parent 761c53c commit 89ae917

13 files changed

+283
-106
lines changed

‎SYCL/Reduction/reduction_big_data.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,7 @@ int test(queue &Q, T Identity) {
5050
// Initialize.
5151
BinaryOperation BOp;
5252
T CorrectOut;
53-
initInputData(InBuf, CorrectOut, Identity, BOp, NWorkItems);
53+
initInputData(InBuf, CorrectOut, BOp, NWorkItems);
5454

5555
// Compute.
5656
Q.submit([&](handler &CGH) {

‎SYCL/Reduction/reduction_ctor.cpp

Lines changed: 11 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
// RUN: %GPU_RUN_PLACEHOLDER %t.out
44
// RUN: %ACC_RUN_PLACEHOLDER %t.out
55

6-
// This performs basic checks such as reduction creation, getIdentity() method,
6+
// This performs basic checks such as reduction creation, identity methods,
77
// and the combine() method of the aux class 'reducer'.
88
// Note: This test relies on non-standard implementation details.
99

@@ -27,8 +27,9 @@ void test_reducer(Reduction &Redu, T A, T B) {
2727
T ExpectedValue = BOp(A, B);
2828
assert(ExpectedValue == detail::ReducerAccess{Reducer}.getElement(0) &&
2929
"Wrong result of binary operation.");
30-
assert(toBool(Reducer.identity() == Redu.getIdentity()) &&
31-
"Failed identity() check().");
30+
assert(
31+
toBool(Reducer.identity() == Redu.getIdentityContainer().getIdentity()) &&
32+
"Failed identity() check().");
3233
}
3334

3435
template <typename T, typename Reduction, typename BinaryOperation>
@@ -41,8 +42,9 @@ void test_reducer(Reduction &Redu, T Identity, BinaryOperation BOp, T A, T B) {
4142
assert(
4243
toBool(ExpectedValue == detail::ReducerAccess{Reducer}.getElement(0)) &&
4344
"Wrong result of binary operation.");
44-
assert(toBool(Reducer.identity() == Redu.getIdentity()) &&
45-
"Failed identity() check().");
45+
assert(
46+
toBool(Reducer.identity() == Redu.getIdentityContainer().getIdentity()) &&
47+
"Failed identity() check().");
4648
}
4749

4850
template <typename... Ts> class KernelNameGroup;
@@ -65,8 +67,8 @@ void testKnown(T Identity, BinaryOperation BOp, T A, T B) {
6567
auto Redu = sycl::reduction(ReduBuf, CGH, BOp);
6668
auto ReduUSM = sycl::reduction(ReduUSMPtr, BOp);
6769

68-
assert(toBool(Redu.getIdentity() == Identity) &&
69-
toBool(ReduUSM.getIdentity() == Identity) &&
70+
assert(toBool(Redu.getIdentityContainer().getIdentity() == Identity) &&
71+
toBool(ReduUSM.getIdentityContainer().getIdentity() == Identity) &&
7072
toBool(known_identity<BinaryOperation, T>::value == Identity) &&
7173
"Failed getIdentity() check().");
7274
test_reducer(Redu, A, B);
@@ -96,8 +98,8 @@ void testUnknown(T Identity, BinaryOperation BOp, T A, T B) {
9698
ReduDWAcc(ReduBuf, CGH);
9799
auto Redu = sycl::reduction(ReduBuf, CGH, Identity, BOp);
98100
auto ReduUSM = sycl::reduction(ReduUSMPtr, Identity, BOp);
99-
assert(toBool(Redu.getIdentity() == Identity) &&
100-
toBool(ReduUSM.getIdentity() == Identity) &&
101+
assert(toBool(Redu.getIdentityContainer().getIdentity() == Identity) &&
102+
toBool(ReduUSM.getIdentityContainer().getIdentity() == Identity) &&
101103
"Failed getIdentity() check().");
102104
test_reducer(Redu, Identity, BOp, A, B);
103105
test_reducer(ReduUSM, Identity, BOp, A, B);

‎SYCL/Reduction/reduction_nd_N_vars.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -48,7 +48,7 @@ struct Red {
4848
}
4949

5050
void init() {
51-
initInputData(InBuf, CorrectOut, IdentityVal, BOp, NWorkItems);
51+
initInputData(InBuf, CorrectOut, BOp, NWorkItems);
5252
if (!PropList.template has_property<
5353
property::reduction::initialize_to_identity>())
5454
CorrectOut = BOp(CorrectOut, InitVal);

‎SYCL/Reduction/reduction_nd_rw.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,13 @@ void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, size_t WGSize,
2222
NumErrors += test<Name>(Q, Identity, Init, BOp, NDRange);
2323
}
2424

25+
template <typename Name, typename T, class BinaryOperation>
26+
void tests(queue &Q, T Init, BinaryOperation BOp, size_t WGSize,
27+
size_t NWItems) {
28+
nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
29+
NumErrors += test<Name>(Q, Init, BOp, NDRange);
30+
}
31+
2532
int main() {
2633
queue Q;
2734
printDeviceInfo(Q);
@@ -52,6 +59,17 @@ int main() {
5259
// Check with CUSTOM type.
5360
using CV = CustomVec<long long>;
5461
tests<class D1, CV>(Q, CV(0), CV(-199), CustomVecPlus<long long>{}, 8, 256);
62+
tests<class D2, CV>(Q, CV(-199), CustomVecPlus<long long>{}, 8, 256);
63+
64+
// Check non power-of-two work-group sizes without identity.
65+
tests<class E1, int>(Q, 99, PlusWithoutIdentity<int>{}, 1, 7);
66+
tests<class E2, int>(Q, -99, PlusWithoutIdentity<int>{}, 49, 49 * 5);
67+
68+
// Try some power-of-two work-group sizes without identity.
69+
tests<class F1, int>(Q, 99, PlusWithoutIdentity<int>{}, 2, 32);
70+
tests<class F2, int>(Q, 199, PlusWithoutIdentity<int>{}, 32, 32);
71+
tests<class F3, int>(Q, 299, PlusWithoutIdentity<int>{}, 128, 256);
72+
tests<class F4, int>(Q, 399, PlusWithoutIdentity<int>{}, 256, 256);
5573

5674
printFinalStatus(NumErrors);
5775
return NumErrors;

‎SYCL/Reduction/reduction_range_1d_rw.cpp

Lines changed: 36 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,11 @@ using namespace sycl;
1212

1313
int NumErrors = 0;
1414

15+
template <typename Name, typename T, typename... ArgTys>
16+
void tests(ArgTys &&...Args) {
17+
NumErrors += test<Name, T>(std::forward<ArgTys>(Args)...);
18+
}
19+
1520
int main() {
1621
queue Q;
1722
printDeviceInfo(Q);
@@ -20,29 +25,41 @@ int main() {
2025

2126
constexpr access::mode RW = access::mode::read_write;
2227
// Fast-reduce and Fast-atomics. Try various range types/sizes.
23-
test<class A1, int>(Q, 0, 99, std::plus<int>{}, range<1>(1));
24-
test<class A2, int>(Q, 0, 99, std::plus<>{}, range<1>(2));
25-
test<class A3, int>(Q, 0, 99, std::plus<>{}, range<1>(7));
26-
test<class A4, int>(Q, 0, 99, std::plus<>{}, range<1>(64));
27-
test<class A5, int>(Q, 0, 99, std::plus<>{}, range<1>(MaxWGSize * 2));
28-
test<class A6, int>(Q, 0, 99, std::plus<>{}, range<1>(MaxWGSize * 2 + 5));
28+
tests<class A1, int>(Q, 0, 99, std::plus<int>{}, range<1>(1));
29+
tests<class A2, int>(Q, 0, 99, std::plus<>{}, range<1>(2));
30+
tests<class A3, int>(Q, 0, 99, std::plus<>{}, range<1>(7));
31+
tests<class A4, int>(Q, 0, 99, std::plus<>{}, range<1>(64));
32+
tests<class A5, int>(Q, 0, 99, std::plus<>{}, range<1>(MaxWGSize * 2));
33+
tests<class A6, int>(Q, 0, 99, std::plus<>{}, range<1>(MaxWGSize * 2 + 5));
2934

3035
// Try various types & ranges.
31-
test<class B1, int>(Q, ~0, ~0, std::bit_and<>{}, range<1>(8));
32-
test<class B2, int>(Q, 0, 0x12340000, std::bit_xor<>{}, range<1>(16));
33-
test<class B3, int>(Q, 0, 0x3400, std::bit_or<>{}, range<1>(MaxWGSize * 3));
34-
test<class B4, uint64_t>(Q, 1, 2, std::multiplies<>{}, range<1>(16));
35-
test<class B5, float>(Q, 1, 3, std::multiplies<>{}, range<1>(11));
36-
test<class B6, int>(Q, (std::numeric_limits<int>::max)(), -99,
37-
ext::oneapi::minimum<>{}, range<1>(MaxWGSize * 2));
38-
test<class B7, int>(Q, (std::numeric_limits<int>::min)(), 99,
39-
ext::oneapi::maximum<>{}, range<1>(8));
36+
tests<class B1, int>(Q, ~0, ~0, std::bit_and<>{}, range<1>(8));
37+
tests<class B2, int>(Q, 0, 0x12340000, std::bit_xor<>{}, range<1>(16));
38+
tests<class B3, int>(Q, 0, 0x3400, std::bit_or<>{}, range<1>(MaxWGSize * 3));
39+
tests<class B4, uint64_t>(Q, 1, 2, std::multiplies<>{}, range<1>(16));
40+
tests<class B5, float>(Q, 1, 3, std::multiplies<>{}, range<1>(11));
41+
tests<class B6, int>(Q, (std::numeric_limits<int>::max)(), -99,
42+
ext::oneapi::minimum<>{}, range<1>(MaxWGSize * 2));
43+
tests<class B7, int>(Q, (std::numeric_limits<int>::min)(), 99,
44+
ext::oneapi::maximum<>{}, range<1>(8));
4045

4146
// Check with CUSTOM type.
42-
using CV = CustomVec<long long>;
43-
test<class C1>(Q, CV(0), CV(99), CustomVecPlus<long long>{}, range<1>(256));
44-
test<class C2>(Q, CV(0), CV(99), CustomVecPlus<long long>{},
45-
range<1>(MaxWGSize * 3));
47+
tests<class C1, CustomVec<long long>>(Q, 0, 99, CustomVecPlus<long long>{},
48+
range<1>(256));
49+
tests<class C2, CustomVec<long long>>(Q, 0, 99, CustomVecPlus<long long>{},
50+
range<1>(MaxWGSize * 3));
51+
tests<class C3, CustomVec<long long>>(Q, 99, CustomVecPlus<long long>{},
52+
range<1>(72));
53+
54+
// Check with identityless operations.
55+
tests<class D1, int>(Q, 99, PlusWithoutIdentity<int>{}, range<1>(1));
56+
tests<class D2, int>(Q, 99, PlusWithoutIdentity<int>{}, range<1>(2));
57+
tests<class D3, int>(Q, 99, PlusWithoutIdentity<int>{}, range<1>(7));
58+
tests<class D4, int>(Q, 99, PlusWithoutIdentity<int>{}, range<1>(64));
59+
tests<class D5, int>(Q, 99, PlusWithoutIdentity<int>{},
60+
range<1>(MaxWGSize * 2));
61+
tests<class D6, int>(Q, 99, PlusWithoutIdentity<int>{},
62+
range<1>(MaxWGSize * 2 + 5));
4663

4764
printFinalStatus(NumErrors);
4865
return NumErrors;

‎SYCL/Reduction/reduction_range_2d_rw.cpp

Lines changed: 21 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -15,9 +15,9 @@ using namespace sycl;
1515

1616
int NumErrors = 0;
1717

18-
template <typename Name, typename T, class BinaryOperation>
19-
void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, range<2> Range) {
20-
NumErrors += test<Name>(Q, Identity, Init, BOp, Range);
18+
template <typename Name, typename T, typename... ArgTys>
19+
void tests(ArgTys &&...Args) {
20+
NumErrors += test<Name, T>(std::forward<ArgTys>(Args)...);
2121
}
2222

2323
int main() {
@@ -46,8 +46,24 @@ int main() {
4646
ext::oneapi::maximum<>{}, range<2>{3, 3});
4747
tests<class B8, float>(Q, 1, 99, std::multiplies<>{}, range<2>{3, 3});
4848

49-
tests<class C1>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
50-
CustomVecPlus<long long>{}, range<2>{33, MaxWGSize});
49+
tests<class C1, CustomVec<long long>>(Q, 0, 99, CustomVecPlus<long long>{},
50+
range<2>{33, MaxWGSize});
51+
tests<class C2, CustomVec<long long>>(Q, 99, CustomVecPlus<long long>{},
52+
range<2>{33, MaxWGSize});
53+
54+
tests<class D1, int>(Q, 99, PlusWithoutIdentity<int>{}, range<2>{1, 1});
55+
tests<class D2, int>(Q, 99, PlusWithoutIdentity<int>{}, range<2>{2, 2});
56+
tests<class D3, int>(Q, 99, PlusWithoutIdentity<int>{}, range<2>{2, 3});
57+
tests<class D4, int>(Q, 99, PlusWithoutIdentity<int>{},
58+
range<2>{MaxWGSize, 1});
59+
tests<class D5, int64_t>(Q, 99, PlusWithoutIdentity<int64_t>{},
60+
range<2>{1, MaxWGSize});
61+
tests<class D6, int64_t>(Q, 99, PlusWithoutIdentity<int64_t>{},
62+
range<2>{2, MaxWGSize * 2});
63+
tests<class D7, int64_t>(Q, 99, PlusWithoutIdentity<int64_t>{},
64+
range<2>{MaxWGSize * 3, 7});
65+
tests<class D8, int64_t>(Q, 99, PlusWithoutIdentity<int64_t>{},
66+
range<2>{3, MaxWGSize * 3});
5167

5268
printFinalStatus(NumErrors);
5369
return NumErrors;

‎SYCL/Reduction/reduction_range_3d_dw.cpp

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -16,9 +16,9 @@ using namespace sycl;
1616

1717
int NumErrors = 0;
1818

19-
template <typename Name, typename T, class BinaryOperation>
20-
void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, range<3> Range) {
21-
NumErrors += test<Name>(Q, Identity, Init, BOp, Range, init_to_identity());
19+
template <typename Name, typename T, typename... ArgTys>
20+
void tests(ArgTys &&...Args) {
21+
NumErrors += test<Name, T>(std::forward<ArgTys>(Args)..., init_to_identity());
2222
}
2323

2424
int main() {
@@ -59,8 +59,8 @@ int main() {
5959
ext::oneapi::maximum<>{}, range<3>{3, MaxWGSize, 3});
6060
tests<class B8, float>(Q, 1, 99, std::multiplies<>{}, range<3>{3, 3, 5});
6161

62-
tests<class C1>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
63-
CustomVecPlus<long long>{}, range<3>{2, 33, MaxWGSize});
62+
tests<class C1, CustomVec<long long>>(Q, 0, 99, CustomVecPlus<long long>{},
63+
range<3>{2, 33, MaxWGSize});
6464

6565
printFinalStatus(NumErrors);
6666
return NumErrors;

‎SYCL/Reduction/reduction_range_3d_rw.cpp

Lines changed: 27 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -15,9 +15,9 @@ using namespace sycl;
1515

1616
int NumErrors = 0;
1717

18-
template <typename Name, typename T, class BinaryOperation>
19-
void tests(queue &Q, T Identity, T Init, BinaryOperation BOp, range<3> Range) {
20-
NumErrors += test<Name>(Q, Identity, Init, BOp, Range);
18+
template <typename Name, typename T, typename... ArgTys>
19+
void tests(ArgTys &&...Args) {
20+
NumErrors += test<Name, T>(std::forward<ArgTys>(Args)...);
2121
}
2222

2323
int main() {
@@ -58,8 +58,30 @@ int main() {
5858
ext::oneapi::maximum<>{}, range<3>{3, MaxWGSize, 3});
5959
tests<class B8, float>(Q, 1, 99, std::multiplies<>{}, range<3>{3, 3, 4});
6060

61-
tests<class C1>(Q, CustomVec<long long>(0), CustomVec<long long>(99),
62-
CustomVecPlus<long long>{}, range<3>{2, 33, MaxWGSize});
61+
tests<class C1, CustomVec<long long>>(Q, 0, 99, CustomVecPlus<long long>{},
62+
range<3>{2, 33, MaxWGSize});
63+
tests<class C2, CustomVec<long long>>(Q, 99, CustomVecPlus<long long>{},
64+
range<3>{2, 33, MaxWGSize});
65+
66+
tests<class D1, int>(Q, 99, PlusWithoutIdentity<int>{}, range<3>{1, 1, 1});
67+
tests<class D2, int>(Q, 99, PlusWithoutIdentity<int>{}, range<3>{2, 2, 2});
68+
tests<class D3, int>(Q, 99, PlusWithoutIdentity<int>{}, range<3>{2, 3, 4});
69+
70+
/* Temporarily disabled
71+
tests<class D4, int64_t>(Q, 99, PlusWithoutIdentity<int64_t>{},
72+
range<3>{1, 1, MaxWGSize + 1});
73+
tests<class D5, int64_t>(Q, 99, PlusWithoutIdentity<int64_t>{},
74+
range<3>{1, MaxWGSize + 1, 1});
75+
tests<class D6, int64_t>(Q, 99, PlusWithoutIdentity<int64_t>{},
76+
range<3>{MaxWGSize + 1, 1, 1});
77+
*/
78+
79+
tests<class D7, int64_t>(Q, 99, PlusWithoutIdentity<int64_t>{},
80+
range<3>{2, 5, MaxWGSize * 2});
81+
tests<class D8, int64_t>(Q, 99, PlusWithoutIdentity<int64_t>{},
82+
range<3>{3, MaxWGSize * 3, 2});
83+
tests<class D9, int64_t>(Q, 99, PlusWithoutIdentity<int64_t>{},
84+
range<3>{MaxWGSize * 3, 8, 4});
6385

6486
printFinalStatus(NumErrors);
6587
return NumErrors;

‎SYCL/Reduction/reduction_range_N_vars.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -48,7 +48,7 @@ struct Red {
4848
}
4949

5050
void init() {
51-
initInputData(InBuf, CorrectOut, IdentityVal, BOp, NWorkItems);
51+
initInputData(InBuf, CorrectOut, BOp, NWorkItems);
5252
if (!PropList.template has_property<
5353
property::reduction::initialize_to_identity>())
5454
CorrectOut = BOp(CorrectOut, InitVal);

‎SYCL/Reduction/reduction_range_lambda.cpp

Lines changed: 34 additions & 28 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,13 @@
1010

1111
using namespace sycl;
1212

13+
int NumErrors = 0;
14+
15+
template <typename Name, typename T, typename... ArgTys>
16+
void tests(ArgTys &&...Args) {
17+
NumErrors += test<Name, T>(std::forward<ArgTys>(Args)...);
18+
}
19+
1320
int main() {
1421
queue Q;
1522
printDeviceInfo(Q);
@@ -18,34 +25,33 @@ int main() {
1825

1926
auto LambdaSum = [](auto x, auto y) { return (x + y); };
2027

21-
int NumErrors = 0;
22-
23-
NumErrors += test<class A1, int>(Q, 0, 99, LambdaSum, range<1>{7});
24-
NumErrors +=
25-
test<class A2, int>(Q, 0, 99, LambdaSum, range<1>{7}, init_to_identity());
26-
27-
NumErrors +=
28-
test<class A3, int>(Q, 0, 99, LambdaSum, range<1>{MaxWGSize + 1});
29-
NumErrors += test<class A4, int>(Q, 0, 99, LambdaSum, range<1>{MaxWGSize + 1},
30-
init_to_identity());
31-
32-
NumErrors += test<class B1, int>(Q, 0, 99, LambdaSum, range<2>{3, 4});
33-
NumErrors += test<class B2, int>(Q, 0, 99, LambdaSum, range<2>{3, 4},
34-
init_to_identity());
35-
36-
NumErrors +=
37-
test<class B3, int>(Q, 0, 99, LambdaSum, range<2>{3, MaxWGSize + 1});
38-
NumErrors += test<class B4, int>(
39-
Q, 0, 99, LambdaSum, range<2>{3, MaxWGSize + 1}, init_to_identity());
40-
41-
NumErrors += test<class C1, int>(Q, 0, 99, LambdaSum, range<3>{2, 3, 4});
42-
NumErrors += test<class C2, int>(Q, 0, 99, LambdaSum, range<3>{2, 3, 4},
43-
init_to_identity());
44-
45-
NumErrors +=
46-
test<class C3, int>(Q, 0, 99, LambdaSum, range<3>{2, 3, MaxWGSize + 1});
47-
NumErrors += test<class C4, int>(
48-
Q, 0, 99, LambdaSum, range<3>{2, 3, MaxWGSize + 1}, init_to_identity());
28+
tests<class A1, int>(Q, 0, 99, LambdaSum, range<1>{7});
29+
tests<class A2, int>(Q, 0, 99, LambdaSum, range<1>{7}, init_to_identity());
30+
tests<class A3, int>(Q, 99, LambdaSum, range<1>{7});
31+
32+
tests<class A4, int>(Q, 0, 99, LambdaSum, range<1>{MaxWGSize + 1});
33+
tests<class A5, int>(Q, 0, 99, LambdaSum, range<1>{MaxWGSize + 1},
34+
init_to_identity());
35+
tests<class A6, int>(Q, 99, LambdaSum, range<1>{MaxWGSize + 1});
36+
37+
tests<class B1, int>(Q, 0, 99, LambdaSum, range<2>{3, 4});
38+
tests<class B2, int>(Q, 0, 99, LambdaSum, range<2>{3, 4}, init_to_identity());
39+
tests<class B3, int>(Q, 99, LambdaSum, range<2>{3, 4});
40+
41+
tests<class B4, int>(Q, 0, 99, LambdaSum, range<2>{3, MaxWGSize + 1});
42+
tests<class B5, int>(Q, 0, 99, LambdaSum, range<2>{3, MaxWGSize + 1},
43+
init_to_identity());
44+
tests<class B6, int>(Q, 99, LambdaSum, range<2>{3, MaxWGSize + 1});
45+
46+
tests<class C1, int>(Q, 0, 99, LambdaSum, range<3>{2, 3, 4});
47+
tests<class C2, int>(Q, 0, 99, LambdaSum, range<3>{2, 3, 4},
48+
init_to_identity());
49+
tests<class C3, int>(Q, 99, LambdaSum, range<3>{2, 3, 4});
50+
51+
tests<class C4, int>(Q, 0, 99, LambdaSum, range<3>{2, 3, MaxWGSize + 1});
52+
tests<class C5, int>(Q, 0, 99, LambdaSum, range<3>{2, 3, MaxWGSize + 1},
53+
init_to_identity());
54+
tests<class C6, int>(Q, 99, LambdaSum, range<3>{2, 3, MaxWGSize + 1});
4955

5056
printFinalStatus(NumErrors);
5157
return NumErrors;

0 commit comments

Comments
 (0)