Skip to content

Commit f96a788

Browse files
v-klochkovv-klochkov
committed
[SYCL] Add test cases for SYCL2020 reductions + initialize_to_identity property
The corresponding changes in SYCL RT: intel/llvm#3410 intel/llvm#3481 Signed-off-by: Vyacheslav N Klochkov <[email protected]>
1 parent cd95be7 commit f96a788

10 files changed

+419
-308
lines changed

‎SYCL/Reduction/reduction_nd_N_vars.cpp

+40-46
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,3 @@
1-
// TODO: level_zero reports an internal error for this test.
2-
// UNSUPPORTED: level_zero
3-
4-
// TODO: Windows implementation of std::tuple is not trivially copiable and
5-
// thus cannot be passed from HOST to DEVICE. Enable the test on Windows when
6-
// SYCL RT gets new type traits having less strict requirements for objects
7-
// being passed to DEVICE.
8-
// UNSUPPORTED: windows
9-
101
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
112
// RUN: %CPU_RUN_PLACEHOLDER %t.out
123
// RUN: %GPU_RUN_PLACEHOLDER %t.out
@@ -36,25 +27,29 @@ constexpr access::mode RW = access::mode::read_write;
3627
constexpr access::mode DW = access::mode::discard_write;
3728

3829
template <typename T>
39-
bool cherkResultIsExpected(int TestCaseNum, T Expected, T Computed) {
30+
bool cherkResultIsExpected(int TestCaseNum, T Expected, T Computed,
31+
bool IsSYCL2020) {
4032
bool Success;
4133
if (!std::is_floating_point<T>::value)
4234
Success = (Expected == Computed);
4335
else
4436
Success = std::abs((Expected / Computed) - 1) < 0.5;
4537

46-
if (!Success)
47-
std::cout << TestCaseNum << ": Expected value = " << Expected
38+
if (!Success) {
39+
std::cerr << "Is SYCL2020 mode: " << IsSYCL2020 << std::endl;
40+
std::cerr << TestCaseNum << ": Expected value = " << Expected
4841
<< ", Computed value = " << Computed << "\n";
42+
}
4943

5044
return Success;
5145
}
5246

5347
// Returns 0 if the test case passed. Otherwise, some non-zero value.
54-
template <class Name, bool IsSYCL2020Mode, typename T1, access::mode Mode1,
48+
template <class Name, bool IsSYCL2020, typename T1, access::mode Mode1,
5549
typename T2, access::mode Mode2, typename T3, access::mode Mode3,
56-
typename T4, class BinaryOperation1, class BinaryOperation2,
57-
class BinaryOperation3, class BinaryOperation4>
50+
typename T4, access::mode Mode4, class BinaryOperation1,
51+
class BinaryOperation2, class BinaryOperation3,
52+
class BinaryOperation4>
5853
int testOne(T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
5954
T2 IdentityVal2, T2 InitVal2, BinaryOperation2 BOp2,
6055
T3 IdentityVal3, T3 InitVal3, BinaryOperation3 BOp3,
@@ -100,8 +95,9 @@ int testOne(T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
10095
CorrectOut2 = BOp2(CorrectOut2, InitVal2);
10196
if (Mode3 == access::mode::read_write)
10297
CorrectOut3 = BOp3(CorrectOut3, InitVal3);
103-
// 4th reduction is USM and this is read_write.
104-
CorrectOut4 = BOp4(CorrectOut4, InitVal4);
98+
// discard_write mode for USM reductions is available only SYCL2020.
99+
if (Mode4 == access::mode::read_write || !IsSYCL2020)
100+
CorrectOut4 = BOp4(CorrectOut4, InitVal4);
105101

106102
// Inititialize data.
107103
{
@@ -123,17 +119,21 @@ int testOne(T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
123119
}
124120

125121
auto NDR = nd_range<1>{range<1>(NWorkItems), range<1>{WGSize}};
126-
if constexpr (IsSYCL2020Mode) {
122+
if constexpr (IsSYCL2020) {
127123
Q.submit([&](handler &CGH) {
128124
auto In1 = InBuf1.template get_access<access::mode::read>(CGH);
129125
auto In2 = InBuf2.template get_access<access::mode::read>(CGH);
130126
auto In3 = InBuf3.template get_access<access::mode::read>(CGH);
131127
auto In4 = InBuf4.template get_access<access::mode::read>(CGH);
132128

133-
auto Redu1 = sycl::reduction(OutBuf1, CGH, IdentityVal1, BOp1);
134-
auto Redu2 = sycl::reduction(OutBuf2, CGH, IdentityVal2, BOp2);
135-
auto Redu3 = sycl::reduction(OutBuf3, CGH, IdentityVal3, BOp3);
136-
auto Redu4 = sycl::reduction(Out4, IdentityVal4, BOp4);
129+
auto Redu1 = sycl::reduction(OutBuf1, CGH, IdentityVal1, BOp1,
130+
getPropertyList<Mode1>());
131+
auto Redu2 = sycl::reduction(OutBuf2, CGH, IdentityVal2, BOp2,
132+
getPropertyList<Mode2>());
133+
auto Redu3 = sycl::reduction(OutBuf3, CGH, IdentityVal3, BOp3,
134+
getPropertyList<Mode3>());
135+
auto Redu4 =
136+
sycl::reduction(Out4, IdentityVal4, BOp4, getPropertyList<Mode4>());
137137

138138
auto Lambda = [=](nd_item<1> NDIt, auto &Sum1, auto &Sum2, auto &Sum3,
139139
auto &Sum4) {
@@ -193,10 +193,10 @@ int testOne(T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
193193
Out4Val = *Out4;
194194
}
195195

196-
Error += cherkResultIsExpected(1, CorrectOut1, Out1[0]) ? 0 : 1;
197-
Error += cherkResultIsExpected(2, CorrectOut2, Out2[0]) ? 0 : 1;
198-
Error += cherkResultIsExpected(3, CorrectOut3, Out3[0]) ? 0 : 1;
199-
Error += cherkResultIsExpected(4, CorrectOut4, Out4Val) ? 0 : 1;
196+
Error += cherkResultIsExpected(1, CorrectOut1, Out1[0], IsSYCL2020) ? 0 : 1;
197+
Error += cherkResultIsExpected(2, CorrectOut2, Out2[0], IsSYCL2020) ? 0 : 1;
198+
Error += cherkResultIsExpected(3, CorrectOut3, Out3[0], IsSYCL2020) ? 0 : 1;
199+
Error += cherkResultIsExpected(4, CorrectOut4, Out4Val, IsSYCL2020) ? 0 : 1;
200200
free(Out4, Q.get_context());
201201
}
202202

@@ -211,45 +211,39 @@ int testOne(T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
211211
// sycl::reduction and sycl::ONEAPI::reduction
212212
template <class Name, typename T1, access::mode Mode1, typename T2,
213213
access::mode Mode2, typename T3, access::mode Mode3, typename T4,
214-
class BinaryOperation1, class BinaryOperation2,
214+
access::mode Mode4, class BinaryOperation1, class BinaryOperation2,
215215
class BinaryOperation3, class BinaryOperation4>
216216
int testBoth(T1 IdentityVal1, T1 InitVal1, BinaryOperation1 BOp1,
217217
T2 IdentityVal2, T2 InitVal2, BinaryOperation2 BOp2,
218218
T3 IdentityVal3, T3 InitVal3, BinaryOperation3 BOp3,
219219
T4 IdentityVal4, T3 InitVal4, BinaryOperation4 BOp4,
220220
usm::alloc AllocType4, size_t NWorkItems, size_t WGSize) {
221221
int Error =
222-
testOne<KName<Name, false>, false, T1, Mode1, T2, Mode2, T3, Mode3, T4>(
223-
IdentityVal1, InitVal1, BOp1, IdentityVal2, InitVal2, BOp2,
224-
IdentityVal3, InitVal3, BOp3, IdentityVal4, InitVal4, BOp4,
225-
AllocType4, NWorkItems, WGSize);
226-
227-
// TODO: property::reduction::initialize_to_identity is not supported yet.
228-
// Thus only read_write mode is tested now.
229-
constexpr access::mode _Mode1 = (Mode1 == DW) ? RW : Mode1;
230-
constexpr access::mode _Mode2 = (Mode2 == DW) ? RW : Mode2;
231-
constexpr access::mode _Mode3 = (Mode3 == DW) ? RW : Mode3;
232-
Error +=
233-
testOne<KName<Name, true>, true, T1, _Mode1, T2, _Mode2, T3, _Mode3, T4>(
234-
IdentityVal1, InitVal1, BOp1, IdentityVal2, InitVal2, BOp2,
235-
IdentityVal3, InitVal3, BOp3, IdentityVal4, InitVal4, BOp4,
236-
AllocType4, NWorkItems, WGSize);
222+
testOne<KName<Name, false>, false, T1, Mode1, T2, Mode2, T3, Mode3, T4,
223+
Mode4>(IdentityVal1, InitVal1, BOp1, IdentityVal2, InitVal2, BOp2,
224+
IdentityVal3, InitVal3, BOp3, IdentityVal4, InitVal4, BOp4,
225+
AllocType4, NWorkItems, WGSize);
226+
227+
Error += testOne<KName<Name, true>, true, T1, Mode1, T2, Mode2, T3, Mode3, T4,
228+
Mode4>(IdentityVal1, InitVal1, BOp1, IdentityVal2, InitVal2,
229+
BOp2, IdentityVal3, InitVal3, BOp3, IdentityVal4,
230+
InitVal4, BOp4, AllocType4, NWorkItems, WGSize);
237231
return Error;
238232
}
239233

240234
int main() {
241-
int Error = testBoth<class FP32Plus16x16, float, DW, int, RW, short, RW, int>(
235+
int Error = testBoth<class Case1, float, DW, int, RW, short, RW, int, RW>(
242236
0, 1000, std::plus<float>{}, 0, 2000, std::plus<>{}, 0, 4000,
243237
std::bit_or<>{}, 0, 8000, std::bit_xor<>{}, usm::alloc::shared, 16, 16);
244238

245239
auto Add = [](auto x, auto y) { return (x + y); };
246-
Error += testBoth<class FP32Plus5x257, float, RW, int, RW, short, DW, int>(
240+
Error += testBoth<class Case2, float, RW, int, RW, short, DW, int, DW>(
247241
0, 1000, std::plus<float>{}, 0, 2000, std::plus<>{}, 0, 4000, Add, 0,
248-
8000, std::bit_xor<int>{}, usm::alloc::device, 5 * (256 + 1), 5);
242+
8000, std::plus<>{}, usm::alloc::device, 5 * (256 + 1), 5);
249243

250244
if (!Error)
251245
std::cout << "Test passed\n";
252246
else
253-
std::cout << Error << " test-cases failed\n";
247+
std::cerr << Error << " test-cases failed\n";
254248
return Error;
255249
}

‎SYCL/Reduction/reduction_nd_ext_double.cpp

+1-1
Original file line numberDiff line numberDiff line change
@@ -11,4 +11,4 @@
1111

1212
#include "reduction_nd_ext_type.hpp"
1313

14-
int main() { return runTests<double>("cl_khr_double"); }
14+
int main() { return runTests<double>("cl_khr_fp64"); }

‎SYCL/Reduction/reduction_nd_ext_type.hpp

+26-21
Original file line numberDiff line numberDiff line change
@@ -14,26 +14,27 @@ constexpr access::mode DW = access::mode::discard_write;
1414

1515
template <typename Name, bool IsSYCL2020Mode, typename T, int Dim,
1616
access::mode Mode, class BinaryOperation>
17-
void test(T Identity, size_t WGSize, size_t NWItems) {
17+
void test(T Identity, T Init, size_t WGSize, size_t NWItems) {
1818
buffer<T, 1> InBuf(NWItems);
1919
buffer<T, 1> OutBuf(1);
2020

2121
// Initialize.
2222
BinaryOperation BOp;
2323
T CorrectOut;
2424
initInputData(InBuf, CorrectOut, Identity, BOp, NWItems);
25-
2625
if (Mode == access::mode::read_write)
27-
(OutBuf.template get_access<access::mode::write>())[0] = Identity;
26+
CorrectOut = BOp(CorrectOut, Init);
27+
28+
(OutBuf.template get_access<access::mode::write>())[0] = Init;
2829

2930
// Compute.
3031
queue Q;
3132
nd_range<1> NDRange(range<1>{NWItems}, range<1>{WGSize});
3233
if constexpr (IsSYCL2020Mode) {
3334
Q.submit([&](handler &CGH) {
3435
auto In = InBuf.template get_access<access::mode::read>(CGH);
35-
auto Redu = sycl::reduction(OutBuf, CGH, Identity, BOp);
36-
36+
auto Redu =
37+
sycl::reduction(OutBuf, CGH, Identity, BOp, getPropertyList<Mode>());
3738
CGH.parallel_for<Name>(NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
3839
Sum.combine(In[NDIt.get_global_linear_id()]);
3940
});
@@ -60,21 +61,20 @@ void test(T Identity, size_t WGSize, size_t NWItems) {
6061
std::cout << "Computed value: " << ComputedOut
6162
<< ", Expected value: " << CorrectOut << ", MaxDiff = " << MaxDiff
6263
<< "\n";
64+
if (IsSYCL2020Mode)
65+
std::cout << std::endl;
6366
assert(0 && "Wrong value.");
6467
}
6568
}
6669

6770
template <typename Name, typename T, int Dim, access::mode Mode,
6871
class BinaryOperation>
69-
void testBoth(T Identity, size_t WGSize, size_t NWItems) {
72+
void testBoth(T Identity, T Init, size_t WGSize, size_t NWItems) {
7073
test<KName<Name, false>, false, T, Dim, Mode, BinaryOperation>(
71-
Identity, WGSize, NWItems);
74+
Identity, Init, WGSize, NWItems);
7275

73-
// TODO: property::reduction::initialize_to_identity is not supported yet.
74-
// Thus only read_write mode is tested now.
75-
constexpr access::mode _Mode = (Mode == DW) ? RW : Mode;
76-
test<KName<Name, true>, true, T, Dim, _Mode, BinaryOperation>(
77-
Identity, WGSize, NWItems);
76+
test<KName<Name, true>, true, T, Dim, Mode, BinaryOperation>(Identity, Init,
77+
WGSize, NWItems);
7878
}
7979

8080
template <typename T> int runTests(const string_class &ExtensionName) {
@@ -84,17 +84,22 @@ template <typename T> int runTests(const string_class &ExtensionName) {
8484
return 0;
8585
}
8686

87-
// Check some less standards WG sizes and corner cases first.
88-
testBoth<class A, T, 1, RW, std::multiplies<T>>(0, 4, 4);
89-
testBoth<class B, T, 0, DW, ONEAPI::plus<T>>(0, 4, 64);
87+
testBoth<class A, T, 1, RW, std::multiplies<T>>(1, 77, 4, 4);
88+
89+
testBoth<class B1, T, 0, DW, ONEAPI::plus<T>>(0, 77, 4, 64);
90+
testBoth<class B2, T, 1, RW, ONEAPI::plus<>>(0, 33, 3, 3 * 5);
9091

91-
testBoth<class C, T, 0, RW, ONEAPI::minimum<T>>(getMaximumFPValue<T>(), 7, 7);
92-
testBoth<class D, T, 1, access::mode::discard_write, ONEAPI::maximum<T>>(
93-
getMinimumFPValue<T>(), 7, 7 * 5);
92+
testBoth<class C1, T, 0, RW, ONEAPI::minimum<T>>(getMaximumFPValue<T>(),
93+
-10.0, 7, 7);
94+
testBoth<class C2, T, 0, RW, ONEAPI::minimum<T>>(getMaximumFPValue<T>(), 99.0,
95+
7, 7);
96+
testBoth<class C3, T, 1, DW, ONEAPI::minimum<>>(getMaximumFPValue<T>(), -99.0,
97+
3, 3);
9498

95-
testBoth<class E, T, 1, RW, ONEAPI::plus<>>(1, 3, 3 * 5);
96-
testBoth<class F, T, 1, DW, ONEAPI::minimum<>>(getMaximumFPValue<T>(), 3, 3);
97-
testBoth<class G, T, 0, DW, ONEAPI::maximum<>>(getMinimumFPValue<T>(), 3, 3);
99+
testBoth<class D1, T, 0, DW, ONEAPI::maximum<>>(getMinimumFPValue<T>(), 99.0,
100+
3, 3);
101+
testBoth<class D2, T, 1, RW, ONEAPI::maximum<T>>(getMinimumFPValue<T>(), 99.0,
102+
7, 7 * 5);
98103

99104
std::cout << "Test passed\n";
100105
return 0;

0 commit comments

Comments
 (0)