|
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