Skip to content

Commit e28efb4

Browse files
haonanya1haonanya1
committed
[SYCL] [Test] Put all math function calls into a single kernel
Signed-off-by: haonanya <[email protected]>
1 parent 3836b41 commit e28efb4

File tree

4 files changed

+104
-224
lines changed

4 files changed

+104
-224
lines changed

‎sycl/test/devicelib/cmath_fp64_test.cpp

Lines changed: 26 additions & 56 deletions
Original file line numberDiff line numberDiff line change
@@ -10,62 +10,29 @@ namespace s = cl::sycl;
1010
constexpr s::access::mode sycl_read = s::access::mode::read;
1111
constexpr s::access::mode sycl_write = s::access::mode::write;
1212

13-
double a = 0;
14-
double b = 1;
15-
double c = 0.5;
16-
double d = 2;
17-
double e = 5;
13+
#define TEST_NUM 38
14+
15+
double ref[TEST_NUM] = {
16+
1, 0, 0, 0, 0, 0, 0, 1, 1, 0.5,
17+
0, 2, 0, 0, 1, 0, 2, 0, 0, 0,
18+
0, 0, 1, 0, 1, 2, 0, 1, 2, 5,
19+
0, 0, 0, 0, 0.5, 0.5, NAN, NAN,};
20+
21+
double refIptr = 1;
1822

1923
template <class T>
2024
void device_cmath_test(s::queue &deviceQueue) {
21-
s::range<1> numOfItems{38};
22-
T result[38] = {-1};
23-
T ref[38] = {
24-
b,
25-
a,
26-
a,
27-
a,
28-
a,
29-
a,
30-
a,
31-
b,
32-
b,
33-
c,
34-
a,
35-
d,
36-
a,
37-
a,
38-
b,
39-
a,
40-
d,
41-
a,
42-
a,
43-
a,
44-
a,
45-
a,
46-
b,
47-
a,
48-
b,
49-
d,
50-
a,
51-
b,
52-
d,
53-
e,
54-
a,
55-
a,
56-
a,
57-
a,
58-
c,
59-
c,
60-
a,
61-
a,
62-
};
63-
int expv = -1;
25+
s::range<1> numOfItems{TEST_NUM};
26+
T result[TEST_NUM] = {-1};
27+
// Variable exponent is an integer value to store the exponent in frexp function
28+
int exponent = -1;
29+
// Variable iptr stores the integral part of float point in modf function
6430
T iptr = -1;
31+
// Variable quo stores the sign and some bits of x/y in remquo function
6532
int quo = -1;
6633
{
6734
s::buffer<T, 1> buffer1(result, numOfItems);
68-
s::buffer<int, 1> buffer2(&expv, s::range<1>{1});
35+
s::buffer<int, 1> buffer2(&exponent, s::range<1>{1});
6936
s::buffer<T, 1> buffer3(&iptr, s::range<1>{1});
7037
s::buffer<int, 1> buffer4(&quo, s::range<1>{1});
7138
deviceQueue.submit([&](cl::sycl::handler &cgh) {
@@ -187,8 +154,9 @@ void device_cmath_test(s::queue &deviceQueue) {
187154
res_access[i++] = std::expm1(a);
188155
}
189156
{
190-
T a = 0;
191-
res_access[i++] = std::fdim(1, a);
157+
T a = 1;
158+
T b = 0;
159+
res_access[i++] = std::fdim(a, b);
192160
}
193161
{
194162
T a = 1;
@@ -238,13 +206,15 @@ void device_cmath_test(s::queue &deviceQueue) {
238206
});
239207
});
240208
}
241-
for (int i = 0; i < 36; ++i) {
209+
// Compare result with reference
210+
for (int i = 0; i < TEST_NUM; ++i) {
242211
assert(is_about_FP(result[i], ref[i]));
243212
}
244-
assert(std::isnan(result[36]));
245-
assert(std::isnan(result[37]));
246-
assert(is_about_FP(iptr, b));
247-
assert(expv == 0);
213+
// Test modf integral part
214+
assert(is_about_FP(iptr, refIptr));
215+
// Test frexp exponent
216+
assert(exponent == 0);
217+
// Test remquo sign
248218
assert(quo == 0);
249219
}
250220

‎sycl/test/devicelib/cmath_test.cpp

Lines changed: 26 additions & 56 deletions
Original file line numberDiff line numberDiff line change
@@ -10,62 +10,29 @@ namespace s = cl::sycl;
1010
constexpr s::access::mode sycl_read = s::access::mode::read;
1111
constexpr s::access::mode sycl_write = s::access::mode::write;
1212

13-
float a = 0;
14-
float b = 1;
15-
float c = 0.5;
16-
float d = 2;
17-
float e = 5;
13+
#define TEST_NUM 38
14+
15+
float ref[TEST_NUM] = {
16+
1, 0, 0, 0, 0, 0, 0, 1, 1, 0.5,
17+
0, 2, 0, 0, 1, 0, 2, 0, 0, 0,
18+
0, 0, 1, 0, 1, 2, 0, 1, 2, 5,
19+
0, 0, 0, 0, 0.5, 0.5, NAN, NAN,};
20+
21+
float refIptr = 1;
1822

1923
template <class T>
2024
void device_cmath_test(s::queue &deviceQueue) {
21-
s::range<1> numOfItems{38};
22-
T result[38] = {-1};
23-
T ref[38] = {
24-
b,
25-
a,
26-
a,
27-
a,
28-
a,
29-
a,
30-
a,
31-
b,
32-
b,
33-
c,
34-
a,
35-
d,
36-
a,
37-
a,
38-
b,
39-
a,
40-
d,
41-
a,
42-
a,
43-
a,
44-
a,
45-
a,
46-
b,
47-
a,
48-
b,
49-
d,
50-
a,
51-
b,
52-
d,
53-
e,
54-
a,
55-
a,
56-
a,
57-
a,
58-
c,
59-
c,
60-
a,
61-
a,
62-
};
63-
int expv = -1;
25+
s::range<1> numOfItems{TEST_NUM};
26+
T result[TEST_NUM] = {-1};
27+
// Variable exponent is an integer value to store the exponent in frexp function
28+
int exponent = -1;
29+
// Variable iptr stores the integral part of float point in modf function
6430
T iptr = -1;
31+
// Variable quo stores the sign and some bits of x/y in remquo function
6532
int quo = -1;
6633
{
6734
s::buffer<T, 1> buffer1(result, numOfItems);
68-
s::buffer<int, 1> buffer2(&expv, s::range<1>{1});
35+
s::buffer<int, 1> buffer2(&exponent, s::range<1>{1});
6936
s::buffer<T, 1> buffer3(&iptr, s::range<1>{1});
7037
s::buffer<int, 1> buffer4(&quo, s::range<1>{1});
7138
deviceQueue.submit([&](cl::sycl::handler &cgh) {
@@ -187,8 +154,9 @@ void device_cmath_test(s::queue &deviceQueue) {
187154
res_access[i++] = std::expm1(a);
188155
}
189156
{
190-
T a = 0;
191-
res_access[i++] = std::fdim(1, a);
157+
T a = 1;
158+
T b = 0;
159+
res_access[i++] = std::fdim(a, b);
192160
}
193161
{
194162
T a = 1;
@@ -238,13 +206,15 @@ void device_cmath_test(s::queue &deviceQueue) {
238206
});
239207
});
240208
}
241-
for (int i = 0; i < 36; ++i) {
209+
// Compare result with reference
210+
for (int i = 0; i < TEST_NUM; ++i) {
242211
assert(is_about_FP(result[i], ref[i]));
243212
}
244-
assert(std::isnan(result[36]));
245-
assert(std::isnan(result[37]));
246-
assert(is_about_FP(iptr, b));
247-
assert(expv == 0);
213+
// Test modf integral part
214+
assert(is_about_FP(iptr, refIptr));
215+
// Test frexp exponent
216+
assert(exponent == 0);
217+
// Test remquo sign
248218
assert(quo == 0);
249219
}
250220

‎sycl/test/devicelib/math_fp64_test.cpp

Lines changed: 26 additions & 56 deletions
Original file line numberDiff line numberDiff line change
@@ -10,61 +10,28 @@ namespace s = cl::sycl;
1010
constexpr s::access::mode sycl_read = s::access::mode::read;
1111
constexpr s::access::mode sycl_write = s::access::mode::write;
1212

13-
double a = 0;
14-
double b = 1;
15-
double c = 0.5;
16-
double d = 2;
17-
double e = 5;
13+
#define TEST_NUM 38
14+
15+
double ref[TEST_NUM] = {
16+
1, 0, 0, 0, 0, 0, 0, 1, 1, 0.5,
17+
0, 2, 0, 0, 1, 0, 2, 0, 0, 0,
18+
0, 0, 1, 0, 1, 2, 0, 1, 2, 5,
19+
0, 0, 0, 0, 0.5, 0.5, NAN, NAN,};
20+
21+
double refIptr = 1;
1822

1923
void device_math_test(s::queue &deviceQueue) {
20-
s::range<1> numOfItems{38};
21-
double result[38] = {-1};
22-
double ref[38] = {
23-
b,
24-
a,
25-
a,
26-
a,
27-
a,
28-
a,
29-
a,
30-
b,
31-
b,
32-
c,
33-
a,
34-
d,
35-
a,
36-
a,
37-
b,
38-
a,
39-
d,
40-
a,
41-
a,
42-
a,
43-
a,
44-
a,
45-
b,
46-
a,
47-
b,
48-
d,
49-
a,
50-
b,
51-
d,
52-
e,
53-
a,
54-
a,
55-
a,
56-
a,
57-
c,
58-
c,
59-
a,
60-
a,
61-
};
62-
int expv = -1;
24+
s::range<1> numOfItems{TEST_NUM};
25+
double result[TEST_NUM] = {-1};
26+
// Variable exponent is an integer value to store the exponent in frexp function
27+
int exponent = -1;
28+
// Variable iptr stores the integral part of float point in modf function
6329
double iptr = -1;
30+
// Variable quo stores the sign and some bits of x/y in remquo function
6431
int quo = -1;
6532
{
6633
s::buffer<double, 1> buffer1(result, numOfItems);
67-
s::buffer<int, 1> buffer2(&expv, s::range<1>{1});
34+
s::buffer<int, 1> buffer2(&exponent, s::range<1>{1});
6835
s::buffer<double, 1> buffer3(&iptr, s::range<1>{1});
6936
s::buffer<int, 1> buffer4(&quo, s::range<1>{1});
7037
deviceQueue.submit([&](cl::sycl::handler &cgh) {
@@ -186,8 +153,9 @@ void device_math_test(s::queue &deviceQueue) {
186153
res_access[i++] = expm1(a);
187154
}
188155
{
189-
double a = 0;
190-
res_access[i++] = fdim(1, a);
156+
double a = 1;
157+
double b = 0;
158+
res_access[i++] = fdim(a, b);
191159
}
192160
{
193161
double a = 1;
@@ -237,13 +205,15 @@ void device_math_test(s::queue &deviceQueue) {
237205
});
238206
});
239207
}
240-
for (int i = 0; i < 36; ++i) {
208+
// Compare result with reference
209+
for (int i = 0; i < TEST_NUM; ++i) {
241210
assert(is_about_FP(result[i], ref[i]));
242211
}
243-
assert(std::isnan(result[36]));
244-
assert(std::isnan(result[37]));
245-
assert(is_about_FP(iptr, b));
246-
assert(expv == 0);
212+
// Test modf integral part
213+
assert(is_about_FP(iptr, refIptr));
214+
// Test frexp exponent
215+
assert(exponent == 0);
216+
// Test remquo sign
247217
assert(quo == 0);
248218
}
249219

0 commit comments

Comments
 (0)