llvm
48737ca5 - [SYCL][TEST] Fix math_test_marray_vec failure on Adreno with ULP (#22867)

Commit
35 days ago
[SYCL][TEST] Fix math_test_marray_vec failure on Adreno with ULP (#22867) 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.
Parents
Loading