[SYCL][TEST] Fix math_test_marray_vec failure on Adreno with ULP - #22867
Conversation
ULP
The test was crashing with exit code 0xC0000409 (assertion failure)
because erfc(0) on the Adreno GPU returns a value 2 ULPs away
from the mathematically exact result:
GPU: 0.99999988 (0x3F7FFFFF)
Expected: 1.0 (0x3F800000)
ULP diff: 2
The root cause is that the original checkEqual overloads used exact
equality (== / !=) which does not account for GPU floating-point
rounding. The OpenCL 3.0 spec permits up to 16 ULPs for erfc, so
the GPU result is well within spec.
A new scalar checkEqual(T a, T b, unsigned maxUlps = 0) overload
is added and placed in a new shared header ulp_utils.hpp at the
test-e2e root so it is available to all tests. The two existing
vec<T,3> and vec<T,4> overloads are replaced by a single generic
vec<T,N> overload, and the marray<T,N> overload is updated to accept
and pass through the maxUlps parameter. A static_assert is added to
restrict checkEqual to float, double, and sycl::half.
The ULP distance is computed by:
1. Short-circuiting on a == b to correctly handle -0 vs +0 (IEEE 754
equal but 1 ULP apart in the bit representation).
2. Reinterpreting float/double/half bits as uint32_t/uint64_t/uint16_t
via memcpy.
3. Remapping sign-magnitude to an ordered integer number line so that
integer subtraction gives the correct ULP distance at all
magnitudes and across the positive/negative boundary.
The tolerance was determined empirically by measuring the actual ULP
diff for all 48 subtests on the Adreno GPU:
- 47/48 subtests: 0 ULPs (bit-exact)
- erfc(0): 2 ULPs (only non-exact result)
The maxUlps parameter defaults to 0 (exact equality) so all 47
bit-exact subtests are unchanged. Only the two erfc call sites
(in math_tests_3 and math_tests_4) pass maxUlps = 2, making the
tolerance precisely targeted to the one function that needs it.
|
@KseniyaTikhomirova @bader |
@vishwajithupendra, known CI issues shouldn't block the merge. This should be ready to go as soon as @intel/llvm-reviewers-runtime team approves. I suggest dropping UPSTREAM tag from the PR title. |
|
I have a concern about this change since it relaxes pass criteria unconditionally for all targets this test is used for. That means that if another backend or HW will have a regression in accuracy comparing to the current one - we may miss it due to relaxed conditions. @vishwajithupendra do you use opencl backend only? |
vishwajithupendra
changed the title
We do use OpenCL only, but I don't think it makes sense to complicate (or duplicate) a lot of tests (because there are plenty more) by enforcing floating point equality "just in case". In the end, these are tests to check SYCL functionality, which allows at least some ULP of difference (depending on the backend), e.g. the CUDA backend allows 5 ULPs. These are not tests to check the accuracy of the HW you're running on. |
|
@intel/llvm-gatekeepers please consider merging this. This PR affects only specific group of tests, failure in |
|
known failures, merging |
5 participants
The test was crashing with exit code 0xC0000409 (assertion failure) because erfc(0) on the Adreno GPU returns a value 2 ULPs away from the mathematically exact result:
GPU: 0.99999988 (0x3F7FFFFF)
Expected: 1.0 (0x3F800000)
ULP diff: 2
The root cause is that the original checkEqual overloads used exact equality (== / !=) which does not account for GPU floating-point rounding. The OpenCL 3.0 spec permits up to 16 ULPs for erfc, so the GPU result is well within spec.
A new scalar checkEqual(T a, T b, unsigned maxUlps = 0) overload is added and placed in a new shared header ulp_utils.hpp at the test-e2e root so it is available to all tests. The two existing vec<T,3> and vec<T,4> overloads are replaced by a single generic vec<T,N> overload, and the marray<T,N> overload is updated to accept and pass through the maxUlps parameter. A static_assert is added to restrict checkEqual to float, double, and sycl::half.
The ULP distance is computed by:
The tolerance was determined empirically by measuring the actual ULP diff for all 48 subtests on the Adreno GPU:
The maxUlps parameter defaults to 0 (exact equality) so all 47 bit-exact subtests are unchanged. Only the two erfc call sites (in math_tests_3 and math_tests_4) pass maxUlps = 2, making the tolerance precisely targeted to the one function that needs it.