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

Commit c873dbc

Browse files
JackAKirkJackAKirk
authored
[SYCL] Tests for vec/marray math (#1002)
Tests for marray/vec SYCL math functions from: intel/llvm#6038 Signed-off-by: jack.kirk <[email protected]>
1 parent e46dbcd commit c873dbc

File tree

7 files changed

+558
-161
lines changed

7 files changed

+558
-161
lines changed

‎SYCL/DeviceLib/built-ins/ext_native_math.cpp

Lines changed: 15 additions & 119 deletions
Original file line numberDiff line numberDiff line change
@@ -3,88 +3,36 @@
33
// RUN: %GPU_RUN_PLACEHOLDER %t.out
44
// RUN: %ACC_RUN_PLACEHOLDER %t.out
55

6-
// OpenCL CPU driver does not support cl_khr_fp16 extension for this reason this
7-
// test is compiled with the -fsycl-device-code-split flag
6+
// Tests oneapi extension native tanh math function for sycl::vec and
7+
// sycl::marray float cases.
88

9-
#include <cassert>
10-
#include <sycl/sycl.hpp>
11-
12-
template <typename T> void assert_out_of_bound(T val, T lower, T upper) {
13-
assert(sycl::all(lower < val && val < upper));
14-
}
15-
16-
template <>
17-
void assert_out_of_bound<float>(float val, float lower, float upper) {
18-
assert(lower < val && val < upper);
19-
}
20-
21-
template <>
22-
void assert_out_of_bound<sycl::half>(sycl::half val, sycl::half lower,
23-
sycl::half upper) {
24-
assert(lower < val && val < upper);
25-
}
26-
27-
template <typename T>
28-
void native_tanh_tester(sycl::queue q, T val, T up, T lo) {
29-
T r = val;
30-
31-
#ifdef SYCL_EXT_ONEAPI_NATIVE_MATH
32-
{
33-
sycl::buffer<T, 1> BufR(&r, sycl::range<1>(1));
34-
q.submit([&](sycl::handler &cgh) {
35-
auto AccR = BufR.template get_access<sycl::access::mode::read_write>(cgh);
36-
cgh.single_task([=]() {
37-
AccR[0] = sycl::ext::oneapi::experimental::native::tanh(AccR[0]);
38-
});
39-
});
40-
}
41-
42-
assert_out_of_bound(r, up, lo);
43-
#else
44-
assert(!"SYCL_EXT_ONEAPI_NATIVE_MATH not supported");
45-
#endif
46-
}
47-
48-
template <typename T>
49-
void native_exp2_tester(sycl::queue q, T val, T up, T lo) {
50-
T r = val;
51-
52-
#ifdef SYCL_EXT_ONEAPI_NATIVE_MATH
53-
{
54-
sycl::buffer<T, 1> BufR(&r, sycl::range<1>(1));
55-
q.submit([&](sycl::handler &cgh) {
56-
auto AccR = BufR.template get_access<sycl::access::mode::read_write>(cgh);
57-
cgh.single_task([=]() {
58-
AccR[0] = sycl::ext::oneapi::experimental::native::exp2(AccR[0]);
59-
});
60-
});
61-
}
62-
63-
assert_out_of_bound(r, up, lo);
64-
#else
65-
assert(!"SYCL_EXT_ONEAPI_NATIVE_MATH not supported");
66-
#endif
67-
}
9+
#include "ext_native_math_common.hpp"
6810

6911
int main() {
7012

7113
sycl::queue q;
7214

73-
const double tv[16] = {-2.0, -1.5, -1.0, 0.0, 2.0, 1.5, 1.0, 0.0,
74-
-1.7, 1.7, -1.2, 1.2, -3.0, 3.0, -10.0, 10.0};
75-
const double tl[16] = {-0.97, -0.91, -0.77, -0.1, 0.95, 0.89, 0.75, -0.1,
76-
-0.94, 0.92, -0.84, 0.82, -1.0, 0.98, -1.10, 0.98};
77-
const double tu[16] = {-0.95, -0.89, -0.75, 0.1, 0.97, 0.91, 0.77, 0.1,
78-
-0.92, 0.94, -0.82, 0.84, -0.98, 1.00, -0.98, 1.10};
15+
const float tv[16] = {-2.0, -1.5, -1.0, 0.0, 2.0, 1.5, 1.0, 0.0,
16+
-1.7, 1.7, -1.2, 1.2, -3.0, 3.0, -10.0, 10.0};
17+
const float tl[16] = {-0.97, -0.91, -0.77, -0.1, 0.95, 0.89, 0.75, -0.1,
18+
-0.94, 0.92, -0.84, 0.82, -1.0, 0.98, -1.10, 0.98};
19+
const float tu[16] = {-0.95, -0.89, -0.75, 0.1, 0.97, 0.91, 0.77, 0.1,
20+
-0.92, 0.94, -0.82, 0.84, -0.98, 1.00, -0.98, 1.10};
7921

8022
native_tanh_tester<float>(q, tv[0], tl[0], tu[0]);
8123
native_tanh_tester<sycl::float2>(q, {tv[0], tv[1]}, {tl[0], tl[1]},
8224
{tu[0], tu[1]});
8325
native_tanh_tester<sycl::float3>(
8426
q, {tv[0], tv[1], tv[2]}, {tl[0], tl[1], tl[2]}, {tu[0], tu[1], tu[2]});
27+
8528
native_tanh_tester<sycl::float4>(q, {tv[0], tv[1], tv[2], tv[3]},
8629
{tl[0], tl[1], tl[2], tl[3]},
8730
{tu[0], tu[1], tu[2], tu[3]});
31+
native_tanh_tester<sycl::marray<float, 3>>(
32+
q, {tv[0], tv[1], tv[2]}, {tl[0], tl[1], tl[2]}, {tu[0], tu[1], tu[2]});
33+
native_tanh_tester<sycl::marray<float, 4>>(q, {tv[0], tv[1], tv[2], tv[3]},
34+
{tl[0], tl[1], tl[2], tl[3]},
35+
{tu[0], tu[1], tu[2], tu[3]});
8836
native_tanh_tester<sycl::float8>(
8937
q, {tv[0], tv[1], tv[2], tv[3], tv[4], tv[5], tv[6], tv[7]},
9038
{tl[0], tl[1], tl[2], tl[3], tl[4], tl[5], tl[6], tl[7]},
@@ -98,57 +46,5 @@ int main() {
9846
{tu[0], tu[1], tu[2], tu[3], tu[4], tu[5], tu[6], tu[7], tu[8], tu[9],
9947
tu[10], tu[11], tu[12], tu[13], tu[14], tu[15]});
10048

101-
if (q.get_device().has(sycl::aspect::fp16)) {
102-
103-
native_tanh_tester<sycl::half>(q, tv[0], tl[0], tu[0]);
104-
native_tanh_tester<sycl::half2>(q, {tv[0], tv[1]}, {tl[0], tl[1]},
105-
{tu[0], tu[1]});
106-
native_tanh_tester<sycl::half3>(
107-
q, {tv[0], tv[1], tv[2]}, {tl[0], tl[1], tl[2]}, {tu[0], tu[1], tu[2]});
108-
native_tanh_tester<sycl::half4>(q, {tv[0], tv[1], tv[2], tv[3]},
109-
{tl[0], tl[1], tl[2], tl[3]},
110-
{tu[0], tu[1], tu[2], tu[3]});
111-
native_tanh_tester<sycl::half8>(
112-
q, {tv[0], tv[1], tv[2], tv[3], tv[4], tv[5], tv[6], tv[7]},
113-
{tl[0], tl[1], tl[2], tl[3], tl[4], tl[5], tl[6], tl[7]},
114-
{tu[0], tu[1], tu[2], tu[3], tu[4], tu[5], tu[6], tu[7]});
115-
native_tanh_tester<sycl::half16>(
116-
q,
117-
{tv[0], tv[1], tv[2], tv[3], tv[4], tv[5], tv[6], tv[7], tv[8], tv[9],
118-
tv[10], tv[11], tv[12], tv[13], tv[14], tv[15]},
119-
{tl[0], tl[1], tl[2], tl[3], tl[4], tl[5], tl[6], tl[7], tl[8], tl[9],
120-
tl[10], tl[11], tl[12], tl[13], tl[14], tl[15]},
121-
{tu[0], tu[1], tu[2], tu[3], tu[4], tu[5], tu[6], tu[7], tu[8], tu[9],
122-
tu[10], tu[11], tu[12], tu[13], tu[14], tu[15]});
123-
124-
const double ev[16] = {-2.0, -1.5, -1.0, 0.0, 2.0, 1.5, 1.0, 0.0,
125-
-2.0, -1.5, -1.0, 0.0, 2.0, 1.5, 1.0, 0.0};
126-
const double el[16] = {0.1, 0.34, 0.4, -0.9, 3.9, 2.7, 1.9, -0.9,
127-
0.1, 0.34, 0.4, -0.9, 3.9, 2.7, 1.9, -0.9};
128-
const double eu[16] = {0.3, 0.36, 0.6, 1.1, 4.1, 2.9, 2.1, 1.1,
129-
0.3, 0.36, 0.6, 1.1, 4.1, 2.9, 2.1, 1.1};
130-
131-
native_exp2_tester<sycl::half>(q, ev[0], el[0], eu[0]);
132-
native_exp2_tester<sycl::half2>(q, {ev[0], ev[1]}, {el[0], el[1]},
133-
{eu[0], eu[1]});
134-
native_exp2_tester<sycl::half3>(
135-
q, {ev[0], ev[1], ev[2]}, {el[0], el[1], el[2]}, {eu[0], eu[1], eu[2]});
136-
native_exp2_tester<sycl::half4>(q, {ev[0], ev[1], ev[2], ev[3]},
137-
{el[0], el[1], el[2], el[3]},
138-
{eu[0], eu[1], eu[2], eu[3]});
139-
native_exp2_tester<sycl::half8>(
140-
q, {ev[0], ev[1], ev[2], ev[3], ev[4], ev[5], ev[6], ev[7]},
141-
{el[0], el[1], el[2], el[3], el[4], el[5], el[6], el[7]},
142-
{eu[0], eu[1], eu[2], eu[3], eu[4], eu[5], eu[6], eu[7]});
143-
native_exp2_tester<sycl::half16>(
144-
q,
145-
{ev[0], ev[1], ev[2], ev[3], ev[4], ev[5], ev[6], ev[7], ev[8], ev[9],
146-
ev[10], ev[11], ev[12], ev[13], ev[14], ev[15]},
147-
{el[0], el[1], el[2], el[3], el[4], el[5], el[6], el[7], el[8], el[9],
148-
el[10], el[11], el[12], el[13], el[14], el[15]},
149-
{eu[0], eu[1], eu[2], eu[3], eu[4], eu[5], eu[6], eu[7], eu[8], eu[9],
150-
eu[10], eu[11], eu[12], eu[13], eu[14], eu[15]});
151-
}
152-
15349
return 0;
15450
}

‎SYCL/DeviceLib/built-ins/ext_native_math_common.hpp

Lines changed: 67 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,67 @@
1+
#include <cassert>
2+
#include <sycl/sycl.hpp>
3+
4+
template <typename T, size_t N>
5+
void assert_out_of_bound(sycl::marray<T, N> val, sycl::marray<T, N> lower,
6+
sycl::marray<T, N> upper) {
7+
for (int i = 0; i < N; i++) {
8+
assert(lower[i] < val[i] && val[i] < upper[i]);
9+
}
10+
}
11+
12+
template <typename T> void assert_out_of_bound(T val, T lower, T upper) {
13+
assert(sycl::all(lower < val && val < upper));
14+
}
15+
16+
template <>
17+
void assert_out_of_bound<float>(float val, float lower, float upper) {
18+
assert(lower < val && val < upper);
19+
}
20+
21+
template <>
22+
void assert_out_of_bound<sycl::half>(sycl::half val, sycl::half lower,
23+
sycl::half upper) {
24+
assert(lower < val && val < upper);
25+
}
26+
27+
template <typename T>
28+
void native_tanh_tester(sycl::queue q, T val, T up, T lo) {
29+
T r = val;
30+
31+
#ifdef SYCL_EXT_ONEAPI_NATIVE_MATH
32+
{
33+
sycl::buffer<T, 1> BufR(&r, sycl::range<1>(1));
34+
q.submit([&](sycl::handler &cgh) {
35+
auto AccR = BufR.template get_access<sycl::access::mode::read_write>(cgh);
36+
cgh.single_task([=]() {
37+
AccR[0] = sycl::ext::oneapi::experimental::native::tanh(AccR[0]);
38+
});
39+
});
40+
}
41+
42+
assert_out_of_bound(r, up, lo);
43+
#else
44+
assert(!"SYCL_EXT_ONEAPI_NATIVE_MATH not supported");
45+
#endif
46+
}
47+
48+
template <typename T>
49+
void native_exp2_tester(sycl::queue q, T val, T up, T lo) {
50+
T r = val;
51+
52+
#ifdef SYCL_EXT_ONEAPI_NATIVE_MATH
53+
{
54+
sycl::buffer<T, 1> BufR(&r, sycl::range<1>(1));
55+
q.submit([&](sycl::handler &cgh) {
56+
auto AccR = BufR.template get_access<sycl::access::mode::read_write>(cgh);
57+
cgh.single_task([=]() {
58+
AccR[0] = sycl::ext::oneapi::experimental::native::exp2(AccR[0]);
59+
});
60+
});
61+
}
62+
63+
assert_out_of_bound(r, up, lo);
64+
#else
65+
assert(!"SYCL_EXT_ONEAPI_NATIVE_MATH not supported");
66+
#endif
67+
}

‎SYCL/DeviceLib/built-ins/ext_native_math_fp16.cpp

Lines changed: 93 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,93 @@
1+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -o %t.out
2+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
3+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
4+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
5+
6+
// OpenCL CPU driver does not support cl_khr_fp16 extension for this reason this
7+
// test is compiled with the -fsycl-device-code-split flag
8+
9+
// Tests oneapi extension native math functions for sycl::vec and sycl::marray
10+
// fp16 cases.
11+
12+
#include "ext_native_math_common.hpp"
13+
14+
int main() {
15+
16+
sycl::queue q;
17+
18+
if (!q.get_device().has(sycl::aspect::fp16)) {
19+
std::cout << "skipping fp16 tests: requires fp16 device aspect."
20+
<< std::endl;
21+
return 0;
22+
}
23+
24+
const sycl::half tv[16] = {-2.0, -1.5, -1.0, 0.0, 2.0, 1.5, 1.0, 0.0,
25+
-1.7, 1.7, -1.2, 1.2, -3.0, 3.0, -10.0, 10.0};
26+
const sycl::half tl[16] = {-0.97, -0.91, -0.77, -0.1, 0.95, 0.89,
27+
0.75, -0.1, -0.94, 0.92, -0.84, 0.82,
28+
-1.0, 0.98, -1.10, 0.98};
29+
const sycl::half tu[16] = {-0.95, -0.89, -0.75, 0.1, 0.97, 0.91,
30+
0.77, 0.1, -0.92, 0.94, -0.82, 0.84,
31+
-0.98, 1.00, -0.98, 1.10};
32+
33+
native_tanh_tester<sycl::half>(q, tv[0], tl[0], tu[0]);
34+
native_tanh_tester<sycl::half2>(q, {tv[0], tv[1]}, {tl[0], tl[1]},
35+
{tu[0], tu[1]});
36+
native_tanh_tester<sycl::half3>(q, {tv[0], tv[1], tv[2]},
37+
{tl[0], tl[1], tl[2]}, {tu[0], tu[1], tu[2]});
38+
native_tanh_tester<sycl::marray<sycl::half, 3>>(
39+
q, {tv[0], tv[1], tv[2]}, {tl[0], tl[1], tl[2]}, {tu[0], tu[1], tu[2]});
40+
native_tanh_tester<sycl::half4>(q, {tv[0], tv[1], tv[2], tv[3]},
41+
{tl[0], tl[1], tl[2], tl[3]},
42+
{tu[0], tu[1], tu[2], tu[3]});
43+
native_tanh_tester<sycl::marray<sycl::half, 4>>(
44+
q, {tv[0], tv[1], tv[2], tv[3]}, {tl[0], tl[1], tl[2], tl[3]},
45+
{tu[0], tu[1], tu[2], tu[3]});
46+
native_tanh_tester<sycl::half8>(
47+
q, {tv[0], tv[1], tv[2], tv[3], tv[4], tv[5], tv[6], tv[7]},
48+
{tl[0], tl[1], tl[2], tl[3], tl[4], tl[5], tl[6], tl[7]},
49+
{tu[0], tu[1], tu[2], tu[3], tu[4], tu[5], tu[6], tu[7]});
50+
native_tanh_tester<sycl::half16>(
51+
q,
52+
{tv[0], tv[1], tv[2], tv[3], tv[4], tv[5], tv[6], tv[7], tv[8], tv[9],
53+
tv[10], tv[11], tv[12], tv[13], tv[14], tv[15]},
54+
{tl[0], tl[1], tl[2], tl[3], tl[4], tl[5], tl[6], tl[7], tl[8], tl[9],
55+
tl[10], tl[11], tl[12], tl[13], tl[14], tl[15]},
56+
{tu[0], tu[1], tu[2], tu[3], tu[4], tu[5], tu[6], tu[7], tu[8], tu[9],
57+
tu[10], tu[11], tu[12], tu[13], tu[14], tu[15]});
58+
59+
const sycl::half ev[16] = {-2.0, -1.5, -1.0, 0.0, 2.0, 1.5, 1.0, 0.0,
60+
-2.0, -1.5, -1.0, 0.0, 2.0, 1.5, 1.0, 0.0};
61+
const sycl::half el[16] = {0.1, 0.34, 0.4, -0.9, 3.9, 2.7, 1.9, -0.9,
62+
0.1, 0.34, 0.4, -0.9, 3.9, 2.7, 1.9, -0.9};
63+
const sycl::half eu[16] = {0.3, 0.36, 0.6, 1.1, 4.1, 2.9, 2.1, 1.1,
64+
0.3, 0.36, 0.6, 1.1, 4.1, 2.9, 2.1, 1.1};
65+
66+
native_exp2_tester<sycl::half>(q, ev[0], el[0], eu[0]);
67+
native_exp2_tester<sycl::half2>(q, {ev[0], ev[1]}, {el[0], el[1]},
68+
{eu[0], eu[1]});
69+
native_exp2_tester<sycl::half3>(q, {ev[0], ev[1], ev[2]},
70+
{el[0], el[1], el[2]}, {eu[0], eu[1], eu[2]});
71+
native_exp2_tester<sycl::half4>(q, {ev[0], ev[1], ev[2], ev[3]},
72+
{el[0], el[1], el[2], el[3]},
73+
{eu[0], eu[1], eu[2], eu[3]});
74+
native_exp2_tester<sycl::marray<sycl::half, 3>>(
75+
q, {ev[0], ev[1], ev[2]}, {el[0], el[1], el[2]}, {eu[0], eu[1], eu[2]});
76+
native_exp2_tester<sycl::marray<sycl::half, 4>>(
77+
q, {ev[0], ev[1], ev[2], ev[3]}, {el[0], el[1], el[2], el[3]},
78+
{eu[0], eu[1], eu[2], eu[3]});
79+
native_exp2_tester<sycl::half8>(
80+
q, {ev[0], ev[1], ev[2], ev[3], ev[4], ev[5], ev[6], ev[7]},
81+
{el[0], el[1], el[2], el[3], el[4], el[5], el[6], el[7]},
82+
{eu[0], eu[1], eu[2], eu[3], eu[4], eu[5], eu[6], eu[7]});
83+
native_exp2_tester<sycl::half16>(
84+
q,
85+
{ev[0], ev[1], ev[2], ev[3], ev[4], ev[5], ev[6], ev[7], ev[8], ev[9],
86+
ev[10], ev[11], ev[12], ev[13], ev[14], ev[15]},
87+
{el[0], el[1], el[2], el[3], el[4], el[5], el[6], el[7], el[8], el[9],
88+
el[10], el[11], el[12], el[13], el[14], el[15]},
89+
{eu[0], eu[1], eu[2], eu[3], eu[4], eu[5], eu[6], eu[7], eu[8], eu[9],
90+
eu[10], eu[11], eu[12], eu[13], eu[14], eu[15]});
91+
92+
return 0;
93+
}

0 commit comments

Comments
 (0)