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

Commit d89605a

Browse files
kbobrovskbobrovs
authored
[ESIMD] Add bfloat16 test cases. (#1193)
* [ESIMD] Add bfloat16 test cases - binary ops, unary plus, memory access, taking a view.
1 parent 6393422 commit d89605a

10 files changed

+160
-31
lines changed

‎SYCL/ESIMD/api/bin_and_cmp_ops_heavy.cpp

Lines changed: 60 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -1,11 +1,12 @@
1-
//==--------------- bin_un_cmp_ops_heavy.cpp - DPC++ ESIMD on-device test -==//
1+
//==-------------- bin_and_cmp_ops_heavy.cpp - DPC++ ESIMD on-device test -==//
22
//
33
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
44
// See https://llvm.org/LICENSE.txt for license information.
55
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
66
//
77
//===----------------------------------------------------------------------===//
8-
// REQUIRES: gpu
8+
// Exclude PVC not to run same test cases twice (via the *_pvc.cpp variant).
9+
// REQUIRES: gpu && !gpu-intel-pvc
910
// UNSUPPORTED: cuda || hip
1011
// RUN: %clangxx -fsycl %s -o %t.out
1112
// RUN: %GPU_RUN_PLACEHOLDER %t.out
@@ -29,6 +30,7 @@
2930

3031
using namespace sycl;
3132
using namespace sycl::ext::intel::esimd;
33+
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
3234

3335
template <class T1, class T2, int VL, class OpClass, class Ops> class TestID;
3436

@@ -68,9 +70,11 @@ template <class T1, class T2, int VL, class OpClass,
6870
template <class, class, class> class VerifyF,
6971
template <class, class, class> class InitF, class Ops>
7072
bool test(Ops ops, queue &q, comp_t<T1, T2, OpClass> epsilon = 0) {
73+
using T = comp_t<T1, T2, OpClass>;
7174
// Log test case info
72-
std::cout << "Testing T1=" << typeid(T1).name() << " T2=" << typeid(T2).name()
73-
<< ", VL=" << VL << " ...\n";
75+
std::cout << "Testing T1=" << esimd_test::type_name<T1>()
76+
<< " T2=" << esimd_test::type_name<T2>() << ", VL=" << VL
77+
<< " comp type: " << esimd_test::type_name<T>() << "...\n";
7478
std::cout << "Operations:";
7579
esimd_test::iterate_ops(ops, [=](OpClass op) {
7680
std::cout << " '" << esimd_test::Op2Str(op) << "'";
@@ -83,7 +87,6 @@ bool test(Ops ops, queue &q, comp_t<T1, T2, OpClass> epsilon = 0) {
8387
T2 *B = sycl::malloc_shared<T2>(Size, q);
8488
constexpr int NumOps = (int)Ops::size;
8589
int CSize = NumOps * Size;
86-
using T = comp_t<T1, T2, OpClass>;
8790
// Result array. For each pair of A[i] and B[i] elements it reserves NumOps
8891
// elements to store result of all operations under test applied to the A[i]
8992
// and B[i]
@@ -181,19 +184,19 @@ template <class T1, class T2, class OpClass> struct verify_strict {
181184
bool operator()(T res, T gold, OpClass op) { return res == gold; }
182185
};
183186

184-
#define EQ(x, y, epsilon) \
185-
((x) > (y) ? (x) - (y) <= epsilon : (y) - (x) <= epsilon)
187+
#define EQ(x, gold, epsilon) \
188+
((x == gold) || (std::abs((double)(x - gold) / (double)gold) <= epsilon))
186189

187-
template <class T1, class T2, class OpClass> struct verify_epsilon {
190+
template <class T1, class T2, class OpClass, bool AllOps = false>
191+
struct verify_epsilon {
188192
using T = comp_t<T1, T2, OpClass>;
189-
T epsilon;
190-
verify_epsilon(T epsilon) : epsilon(epsilon) {}
193+
double epsilon;
194+
verify_epsilon(double epsilon) : epsilon(epsilon) {}
191195

192196
bool operator()(T res, T gold, OpClass op) {
193-
if constexpr (std::is_same_v<OpClass, esimd_test::BinaryOp>) {
194-
if (op == esimd_test::BinaryOp::div) {
195-
return EQ(res, gold, epsilon);
196-
}
197+
if (AllOps || ((std::is_same_v<OpClass, esimd_test::BinaryOp>)&&(
198+
op == esimd_test::BinaryOp::div))) {
199+
return EQ(res, gold, epsilon);
197200
}
198201
return res == gold;
199202
}
@@ -245,6 +248,8 @@ template <class T1, class T2, class OpClass> struct init_for_shift {
245248
// shortcuts for less clutter
246249
template <class T1, class T2, class C> using VSf = verify_strict<T1, T2, C>;
247250
template <class T1, class T2, class C> using VEf = verify_epsilon<T1, T2, C>;
251+
template <class T1, class T2, class C>
252+
using VEfa = verify_epsilon<T1, T2, C, true>;
248253
template <class T1, class T2, class C> using VNf = verify_n<T1, T2, C>;
249254
template <class T1, class T2, class C> using IDf = init_default<T1, T2, C>;
250255
template <class T1, class T2, class C> using ISf = init_for_shift<T1, T2, C>;
@@ -257,7 +262,7 @@ int main(void) {
257262
bool passed = true;
258263
using BinOp = esimd_test::BinaryOp;
259264

260-
auto arith_ops = esimd_test::ArithBinaryOps;
265+
auto arith_ops = esimd_test::ArithBinaryOpsNoDiv;
261266
passed &= test<unsigned char, int, 1, BinOp, VSf, IDf>(arith_ops, q);
262267
passed &= test<char, float, 7, BinOp, VEf, IDf>(arith_ops, q, 0.000001f);
263268
passed &= test<short, double, 7, BinOp, VEf, IDf>(arith_ops, q, 1e-15);
@@ -266,16 +271,49 @@ int main(void) {
266271
passed &= test<half, unsigned int, 32, BinOp, VSf, IDf>(arith_ops, q, 1);
267272
passed &= test<double, half, 7, BinOp, VSf, IDf>(arith_ops, q);
268273
passed &= test<short, uint64_t, 7, BinOp, VSf, IDf>(arith_ops, q);
269-
270-
auto int_ops =
271-
esimd_test::IntBinaryOpsNoShift; // different data needed for shift
274+
#ifdef USE_BF16
275+
passed &= test<bfloat16, int, 8, BinOp, VSf, IDf>(arith_ops, q);
276+
passed &= test<half, bfloat16, 7, BinOp, VEfa, IDf>(arith_ops, q, 0.03);
277+
#endif // USE_BF16
278+
279+
// Test division separately, as error probability is higher.
280+
auto div_op = esimd_test::BinaryOpSeq<BinOp::div>{};
281+
passed &= test<unsigned char, int, 1, BinOp, VSf, IDf>(div_op, q);
282+
passed &= test<char, float, 7, BinOp, VEf, IDf>(div_op, q, 0.000001f);
283+
#ifndef WA_BUG
284+
passed &= test<short, double, 7, BinOp, VSf, IDf>(div_op, q);
285+
#endif // WA_BUG
286+
passed &= test<float, float, 32, BinOp, VEf, IDf>(div_op, q, 0.000001f);
287+
passed &= test<half, char, 1, BinOp, verify_n, IDf>(div_op, q, 1);
288+
passed &= test<half, unsigned int, 32, BinOp, VSf, IDf>(div_op, q, 1);
289+
#ifndef WA_BUG
290+
passed &= test<double, half, 7, BinOp, VSf, IDf>(div_op, q);
291+
#endif // WA_BUG
292+
passed &= test<short, uint64_t, 7, BinOp, VSf, IDf>(div_op, q);
293+
#ifdef USE_BF16
294+
passed &= test<bfloat16, short, 8, BinOp, VSf, IDf>(div_op, q);
295+
passed &= test<half, bfloat16, 7, BinOp, VEfa, IDf>(div_op, q, 0.03);
296+
#endif // USE_BF16
297+
298+
auto int_ops = esimd_test::IntBinaryOpsNoShiftNoDivRem;
272299
passed &= test<unsigned char, unsigned int, 1, BinOp, VSf, IDf>(int_ops, q);
273300
passed &= test<char, uint64_t, 1, BinOp, VSf, IDf>(int_ops, q);
274301
passed &= test<uint64_t, char, 32, BinOp, VSf, IDf>(int_ops, q);
275302
passed &= test<int, short, 1, BinOp, VSf, IDf>(int_ops, q);
276303
passed &= test<short, int, 8, BinOp, VSf, IDf>(int_ops, q);
277304
passed &= test<int, int, 7, BinOp, VSf, IDf>(int_ops, q);
278305

306+
auto int_div_ops = esimd_test::IntBinaryOpsDivRem;
307+
passed &=
308+
test<unsigned char, unsigned int, 1, BinOp, VSf, IDf>(int_div_ops, q);
309+
#ifndef WA_BUG
310+
passed &= test<char, uint64_t, 1, BinOp, VSf, IDf>(int_div_ops, q);
311+
#endif // WA_BUG
312+
passed &= test<uint64_t, char, 32, BinOp, VSf, IDf>(int_div_ops, q);
313+
passed &= test<int, short, 1, BinOp, VSf, IDf>(int_div_ops, q);
314+
passed &= test<short, int, 8, BinOp, VSf, IDf>(int_div_ops, q);
315+
passed &= test<int, int, 7, BinOp, VSf, IDf>(int_div_ops, q);
316+
279317
auto sh_ops = esimd_test::BinaryOpSeq<BinOp::shl, BinOp::shr>{};
280318
passed &= test<unsigned char, unsigned int, 1, BinOp, VSf, ISf>(sh_ops, q);
281319
passed &= test<char, int64_t, 1, BinOp, VSf, ISf>(sh_ops, q);
@@ -294,6 +332,10 @@ int main(void) {
294332
passed &= test<half, unsigned int, 32, CmpOp, VSf, IDf>(cmp_ops, q, 1);
295333
passed &= test<double, half, 7, CmpOp, VSf, IDf>(cmp_ops, q);
296334
passed &= test<short, uint64_t, 7, CmpOp, VSf, IDf>(cmp_ops, q);
335+
#ifdef USE_BF16
336+
passed &= test<bfloat16, int, 32, CmpOp, VSf, IDf>(cmp_ops, q);
337+
passed &= test<half, bfloat16, 7, CmpOp, VSf, IDf>(cmp_ops, q);
338+
#endif // USE_BF16
297339

298340
std::cout << (passed ? "Test PASSED\n" : "Test FAILED\n");
299341
return passed ? 0 : 1;

‎SYCL/ESIMD/api/bin_and_cmp_ops_heavy_pvc.cpp

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,20 @@
1+
//==---------- bin_and_cmp_ops_heavy_pvc.cpp - DPC++ ESIMD on-device test -==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu-intel-pvc
9+
// RUN: %clangxx -fsycl %s -o %t.out
10+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
11+
12+
// Tests various binary operations applied to simd objects.
13+
// PVC variant of the test - adds bfloat16.
14+
15+
// TODO Re-enable cases disabled via WA_BUG.
16+
17+
#define USE_BF16
18+
#define WA_BUG
19+
20+
#include "bin_and_cmp_ops_heavy.cpp"

‎SYCL/ESIMD/api/replicate_smoke.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -20,6 +20,7 @@
2020

2121
using namespace sycl;
2222
using namespace sycl::ext::intel::esimd;
23+
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
2324

2425
template <class T> struct char_to_int {
2526
using type = typename std::conditional<
@@ -178,6 +179,7 @@ int main(int argc, char **argv) {
178179
bool passed = true;
179180

180181
passed &= test<half>(q);
182+
passed &= test<bfloat16>(q);
181183
passed &= test<unsigned char>(q);
182184
passed &= test<short>(q);
183185
passed &= test<unsigned short>(q);

‎SYCL/ESIMD/api/simd_copy_to_from.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -38,6 +38,7 @@
3838
using namespace sycl;
3939
using namespace sycl::ext::intel;
4040
using namespace sycl::ext::intel::esimd;
41+
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
4142

4243
template <typename T, int N, typename Flags>
4344
bool testUSM(queue &Q, T *Src, T *Dst, unsigned Off, Flags) {
@@ -250,9 +251,11 @@ int main(void) {
250251
#else
251252
Pass &= testUSM<uint16_t>(Q);
252253
Pass &= testUSM<float>(Q);
254+
Pass &= testUSM<bfloat16>(Q);
253255

254256
Pass &= testAcc<int16_t>(Q);
255257
Pass &= testAcc<float>(Q);
258+
Pass &= testAcc<bfloat16>(Q);
256259
#endif
257260

258261
std::cout << (Pass ? "Test Passed\n" : "Test FAILED\n");

‎SYCL/ESIMD/api/simd_subscript_operator.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -24,6 +24,7 @@
2424

2525
using namespace sycl;
2626
using namespace sycl::ext::intel::esimd;
27+
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
2728

2829
template <class T> bool test(queue &q) {
2930
std::cout << "Testing " << typeid(T).name() << "...\n";
@@ -102,6 +103,7 @@ int main(int argc, char **argv) {
102103
passed &= test<unsigned int>(q);
103104
passed &= test<float>(q);
104105
passed &= test<half>(q);
106+
passed &= test<bfloat16>(q);
105107

106108
return passed ? 0 : 1;
107109
}

‎SYCL/ESIMD/api/simd_view_subscript_operator.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,7 @@
2525

2626
using namespace sycl;
2727
using namespace sycl::ext::intel::esimd;
28+
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
2829

2930
template <class T> class TestID;
3031

@@ -102,6 +103,7 @@ int main(int argc, char **argv) {
102103
bool passed = true;
103104
passed &= test<int>(q);
104105
passed &= test<half>(q);
106+
passed &= test<bfloat16>(q);
105107

106108
std::cout << (passed ? "Test Passed\n" : "Test FAILED\n");
107109

‎SYCL/ESIMD/api/svm_gather_scatter.cpp

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,7 @@
2525
using namespace sycl;
2626
using namespace sycl::ext::intel;
2727
using namespace sycl::ext::intel::esimd;
28+
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
2829

2930
template <typename T, int N> bool test(queue &Q) {
3031
std::cout << " Running " << typeid(T).name() << " test, N=" << N << "...\n";
@@ -106,6 +107,14 @@ int main(void) {
106107
Pass &= test<half, 8>(Q);
107108
Pass &= test<half, 16>(Q);
108109
Pass &= test<half, 32>(Q);
110+
111+
Pass &= test<bfloat16, 1>(Q);
112+
Pass &= test<bfloat16, 2>(Q);
113+
Pass &= test<bfloat16, 4>(Q);
114+
Pass &= test<bfloat16, 8>(Q);
115+
Pass &= test<bfloat16, 16>(Q);
116+
Pass &= test<bfloat16, 32>(Q);
117+
109118
std::cout << (Pass ? "Test Passed\n" : "Test FAILED\n");
110119
return Pass ? 0 : 1;
111120
}

‎SYCL/ESIMD/api/unary_ops_heavy.cpp

Lines changed: 19 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -5,7 +5,8 @@
55
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
66
//
77
//===----------------------------------------------------------------------===//
8-
// REQUIRES: gpu
8+
// Exclude PVC not to run same test cases twice (via the *_pvc.cpp variant).
9+
// REQUIRES: gpu && !gpu-intel-pvc
910
// UNSUPPORTED: cuda || hip
1011
// RUN: %clangxx -fsycl %s -o %t.out
1112
// RUN: %GPU_RUN_PLACEHOLDER %t.out
@@ -29,6 +30,7 @@
2930

3031
using namespace sycl;
3132
using namespace sycl::ext::intel::esimd;
33+
using bfloat16 = sycl::ext::oneapi::experimental::bfloat16;
3234

3335
template <class T, int VL, class Ops> class TestID;
3436

@@ -51,7 +53,8 @@ template <class T, int VL, class Ops, template <class, int> class SimdT = simd>
5153
bool test(Ops ops, queue &q) {
5254
using OpClass = esimd_test::UnaryOp;
5355
// Log test case info
54-
std::cout << "Testing T=" << typeid(T).name() << ", VL=" << VL << " ...\n";
56+
std::cout << "Testing T=" << esimd_test::type_name<T>() << ", VL=" << VL
57+
<< " ...\n";
5558
std::cout << "Operations:";
5659
esimd_test::iterate_ops(ops, [=](OpClass op) {
5760
std::cout << " '" << esimd_test::Op2Str(op) << "'";
@@ -172,14 +175,20 @@ int main(void) {
172175
passed &= test<float, 32>(mod_ops, q);
173176
passed &= test<double, 7>(mod_ops, q);
174177

175-
auto singed_ops = esimd_test::OpSeq<UnOp, UnOp::minus, UnOp::plus>{};
176-
passed &= test<char, 7>(singed_ops, q);
177-
passed &= test<short, 7>(singed_ops, q);
178-
passed &= test<int, 16>(singed_ops, q);
179-
passed &= test<int64_t, 16>(singed_ops, q);
180-
passed &= test<half, 16>(singed_ops, q);
181-
passed &= test<float, 16>(singed_ops, q);
182-
passed &= test<double, 16>(singed_ops, q);
178+
auto signed_ops = esimd_test::OpSeq<UnOp, UnOp::minus, UnOp::plus>{};
179+
passed &= test<char, 7>(signed_ops, q);
180+
passed &= test<short, 7>(signed_ops, q);
181+
passed &= test<int, 16>(signed_ops, q);
182+
passed &= test<int64_t, 16>(signed_ops, q);
183+
passed &= test<half, 16>(signed_ops, q);
184+
passed &= test<float, 16>(signed_ops, q);
185+
passed &= test<double, 16>(signed_ops, q);
186+
187+
#ifdef USE_BF16
188+
// TODO: the rest unary operations are not yet supported for bfloat16 on host.
189+
auto unary_plus_op = esimd_test::OpSeq<UnOp, UnOp::plus>{};
190+
passed &= test<bfloat16, 16>(unary_plus_op, q);
191+
#endif // USE_BF16
183192

184193
auto bit_ops = esimd_test::OpSeq<UnOp, UnOp::bit_not>{};
185194
passed &= test<char, 7>(bit_ops, q);

‎SYCL/ESIMD/api/unary_ops_heavy_pvc.cpp

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,17 @@
1+
//==--------------- unary_ops_heavy_pvc.cpp - DPC++ ESIMD on-device test --==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu-intel-pvc
9+
// RUN: %clangxx -fsycl %s -o %t.out
10+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
11+
12+
// Tests various unary operations applied to simd objects.
13+
// PVC variant of the test - adds bfloat16.
14+
15+
#define USE_BF16
16+
17+
#include "unary_ops_heavy.cpp"

‎SYCL/ESIMD/esimd_test_utils.hpp

Lines changed: 26 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -438,16 +438,21 @@ static constexpr BinaryOpSeq<BinaryOp::add, BinaryOp::sub, BinaryOp::mul,
438438
BinaryOp::div>
439439
ArithBinaryOps{};
440440

441+
static constexpr BinaryOpSeq<BinaryOp::add, BinaryOp::sub, BinaryOp::mul>
442+
ArithBinaryOpsNoDiv{};
443+
441444
static constexpr BinaryOpSeq<BinaryOp::add, BinaryOp::sub, BinaryOp::mul,
442445
BinaryOp::div, BinaryOp::rem, BinaryOp::shl,
443446
BinaryOp::shr, BinaryOp::bit_or, BinaryOp::bit_and,
444447
BinaryOp::bit_xor>
445448
IntBinaryOps{};
446449

447450
static constexpr BinaryOpSeq<BinaryOp::add, BinaryOp::sub, BinaryOp::mul,
448-
BinaryOp::div, BinaryOp::rem, BinaryOp::bit_or,
449-
BinaryOp::bit_and, BinaryOp::bit_xor>
450-
IntBinaryOpsNoShift{};
451+
BinaryOp::bit_or, BinaryOp::bit_and,
452+
BinaryOp::bit_xor>
453+
IntBinaryOpsNoShiftNoDivRem{};
454+
455+
static constexpr BinaryOpSeq<BinaryOp::div, BinaryOp::rem> IntBinaryOpsDivRem{};
451456

452457
static constexpr OpSeq<CmpOp, CmpOp::lt, CmpOp::lte, CmpOp::eq, CmpOp::ne,
453458
CmpOp::gte, CmpOp::gt>
@@ -538,4 +543,22 @@ std::unique_ptr<T, USMDeleter> usm_malloc_shared(queue q, int n) {
538543
return std::move(res);
539544
}
540545

546+
template <class T> static const char *type_name();
547+
#define TID(T) \
548+
template <> const char *type_name<T>() { return #T; }
549+
TID(char) // for some reason, 'char' does not match 'int8_t' during
550+
// 'type_name' specialization
551+
TID(int8_t)
552+
TID(uint8_t)
553+
TID(int16_t)
554+
TID(uint16_t)
555+
TID(int32_t)
556+
TID(uint32_t)
557+
TID(int64_t)
558+
TID(uint64_t)
559+
TID(half)
560+
TID(sycl::ext::oneapi::experimental::bfloat16)
561+
TID(float)
562+
TID(double)
563+
541564
} // namespace esimd_test

0 commit comments

Comments
 (0)