|
11 | 11 | constexpr float value = 560.0f;
|
12 | 12 | constexpr float divider = 280.0f;
|
13 | 13 |
|
14 |
| -// Reference |
15 |
| -// https://github.com/KhronosGroup/SYCL-CTS/blob/SYCL-2020/util/accuracy.h |
16 |
| -template <typename T> T get_ulp_std(T x) { |
17 |
| - const T inf = std::numeric_limits<T>::infinity(); |
18 |
| - const T negative = std::fabs(std::nextafter(x, -inf) - x); |
19 |
| - const T positive = std::fabs(std::nextafter(x, inf) - x); |
20 |
| - return std::fmin(negative, positive); |
21 |
| -} |
| 14 | +int32_t ulp_difference(float lhs, float rhs) { |
| 15 | + int32_t lhsInt = *reinterpret_cast<int32_t *>(&lhs); |
| 16 | + int32_t rhsInt = *reinterpret_cast<int32_t *>(&rhs); |
22 | 17 |
|
23 |
| -template <typename T> int ulp_difference(const T &lhs, const T &rhs) { |
24 |
| - return get_ulp_std(lhs) - get_ulp_std(rhs); |
| 18 | + return std::abs(lhsInt - rhsInt); |
25 | 19 | }
|
26 | 20 |
|
27 | 21 | void test_div() {
|
28 | 22 | sycl::queue q(sycl::default_selector_v);
|
29 |
| - float *in_value = (float *)sycl::malloc_shared(sizeof(float), q); |
30 |
| - float *in_divider = (float *)sycl::malloc_shared(sizeof(float), q); |
| 23 | + float *inValue = (float *)sycl::malloc_shared(sizeof(float), q); |
| 24 | + float *inDivider = (float *)sycl::malloc_shared(sizeof(float), q); |
31 | 25 | float *output = (float *)sycl::malloc_shared(sizeof(float), q);
|
32 |
| - *in_value = value; |
33 |
| - *in_divider = divider; |
| 26 | + *inValue = value; |
| 27 | + *inDivider = divider; |
34 | 28 | q.submit([&](sycl::handler &h) {
|
35 | 29 | h.single_task([=] {
|
36 |
| - float res = *in_value / *in_divider; |
| 30 | + float res = *inValue / *inDivider; |
37 | 31 | *output = res;
|
38 | 32 | });
|
39 | 33 | }).wait();
|
40 | 34 |
|
41 | 35 | float hostRef = value / divider;
|
42 |
| - int ulpDiff = ulp_difference<float>(hostRef, *output); |
| 36 | + int ulpDiff = ulp_difference(hostRef, *output); |
43 | 37 | assert(std::abs(ulpDiff) < 1 && "Division is not precise");
|
44 | 38 | }
|
45 | 39 |
|
46 | 40 | void test_sqrt() {
|
47 | 41 | sycl::queue q(sycl::default_selector_v);
|
48 |
| - float *in_value = (float *)sycl::malloc_shared(sizeof(float), q); |
| 42 | + float *inValue = (float *)sycl::malloc_shared(sizeof(float), q); |
49 | 43 | float *output = (float *)sycl::malloc_shared(sizeof(float), q);
|
50 |
| - *in_value = value; |
| 44 | + *inValue = value; |
51 | 45 | q.submit([&](sycl::handler &h) {
|
52 | 46 | h.single_task([=] {
|
53 |
| - float res = sycl::sqrt(*in_value); |
| 47 | + float res = sycl::sqrt(*inValue); |
54 | 48 | *output = res;
|
55 | 49 | });
|
56 | 50 | }).wait();
|
57 | 51 |
|
58 | 52 | float hostRef = std::sqrt(value);
|
59 |
| - int ulpDiff = ulp_difference<float>(hostRef, *output); |
| 53 | + int ulpDiff = ulp_difference(hostRef, *output); |
60 | 54 | assert(std::abs(ulpDiff) < 1 && "Sqrt is not precise");
|
61 | 55 | }
|
62 | 56 |
|
|
0 commit comments