From df2ad8f7ffa8b499f13bf66abcd6b0d764382d52 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 7 Dec 2023 01:13:48 -0800 Subject: [PATCH 1/3] Do not check for undefined behavior in abs_diff Similar to sycl::abs, sycl::abs_diff now has undefined behavior for unrepresentable results. This patch changes the generated tests for abs_diff to avoid UB, similar to how it was done for abs in https://github.com/KhronosGroup/SYCL-CTS/pull/831. Signed-off-by: Larsen, Steffen --- util/math_reference.h | 40 ++++++++++++++++++++++++++++++++-------- 1 file changed, 32 insertions(+), 8 deletions(-) diff --git a/util/math_reference.h b/util/math_reference.h index a5782db3d..cc9ad67be 100644 --- a/util/math_reference.h +++ b/util/math_reference.h @@ -305,21 +305,45 @@ sycl_cts::resultRef> abs(sycl::marray a) { /* absolute difference */ template -T abs_diff(T a, T b) { - T h = (a > b) ? a : b; - T l = (a <= b) ? a : b; - return h - l; +sycl_cts::resultRef abs_diff(T a, T b) { + if constexpr (std::is_unsigned_v) { + T h = (a > b) ? a : b; + T l = (a <= b) ? a : b; + return h - l; + } else { + using U = std::make_unsigned_t; + constexpr T TMax = std::numeric_limits::max(); + // For unsigned integers, negative values might not be representable by + // an absolute value, so we compute the overflow separately. + T ao = a < -TMax ? a + TMax : 0; + T bo = b < -TMax ? b + TMax : 0; + // Then find the absolute values (without the overflow) and compute the + // unsigned difference. + U au = static_cast(std::abs(a - ao)); + U bu = static_cast(std::abs(b - bo)); + U h = (au > bu) ? au : bu; + U l = (au <= bu) ? au : bu; + U result = (a < 0) == (b < 0) ? h - l : h + l; + // Now re-add the absolute distance in overflow. Both ao and bo will be + // <= 0. + result += (ao <= bo ? U(bo - ao) : U(ao - bo)); + // If the unsigned difference is larger than the signed maximum value, we + // have UB. Otherwise we have our valid result. + return result > U(TMax) ? sycl_cts::resultRef(0, true) : T(result); + } } template -sycl::vec abs_diff(sycl::vec a, sycl::vec b) { - return sycl_cts::math::run_func_on_vector( +sycl_cts::resultRef> abs_diff(sycl::vec a, + sycl::vec b) { + return sycl_cts::math::run_func_on_vector_result_ref( [](T x, T y) { return abs_diff(x, y); }, a, b); } // FIXME: hipSYCL does not support marray #ifndef SYCL_CTS_COMPILING_WITH_HIPSYCL template -sycl::marray abs_diff(sycl::marray a, sycl::marray b) { - return sycl_cts::math::run_func_on_marray( +sycl_cts::resultRef> abs_diff(sycl::marray a, + sycl::marray b) { + return sycl_cts::math::run_func_on_marray_result_ref( [](T x, T y) { return abs_diff(x, y); }, a, b); } #endif From 51fdb958c5e7e135f4a9305c442a21176103a25d Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Fri, 15 Dec 2023 08:44:59 +0100 Subject: [PATCH 2/3] Update util/math_reference.h --- util/math_reference.h | 34 +++++++++------------------------- 1 file changed, 9 insertions(+), 25 deletions(-) diff --git a/util/math_reference.h b/util/math_reference.h index cc9ad67be..533025aef 100644 --- a/util/math_reference.h +++ b/util/math_reference.h @@ -306,31 +306,15 @@ sycl_cts::resultRef> abs(sycl::marray a) { /* absolute difference */ template sycl_cts::resultRef abs_diff(T a, T b) { - if constexpr (std::is_unsigned_v) { - T h = (a > b) ? a : b; - T l = (a <= b) ? a : b; - return h - l; - } else { - using U = std::make_unsigned_t; - constexpr T TMax = std::numeric_limits::max(); - // For unsigned integers, negative values might not be representable by - // an absolute value, so we compute the overflow separately. - T ao = a < -TMax ? a + TMax : 0; - T bo = b < -TMax ? b + TMax : 0; - // Then find the absolute values (without the overflow) and compute the - // unsigned difference. - U au = static_cast(std::abs(a - ao)); - U bu = static_cast(std::abs(b - bo)); - U h = (au > bu) ? au : bu; - U l = (au <= bu) ? au : bu; - U result = (a < 0) == (b < 0) ? h - l : h + l; - // Now re-add the absolute distance in overflow. Both ao and bo will be - // <= 0. - result += (ao <= bo ? U(bo - ao) : U(ao - bo)); - // If the unsigned difference is larger than the signed maximum value, we - // have UB. Otherwise we have our valid result. - return result > U(TMax) ? sycl_cts::resultRef(0, true) : T(result); - } + using U = std::make_unsigned_t; + T h = (a > b) ? a : b; + T l = (a <= b) ? a : b; + // Use knowledge of 2-complement encoding while we know the difference is + // positive, so no overflow/underflow. + U result = static_cast(h) - static_cast(l); + return result > static_cast(std::numeric_limits::max()) + ? sycl_cts::resultRef(0, true) + : T(result); } template sycl_cts::resultRef> abs_diff(sycl::vec a, From 860addeba9b4752e3ea71c860842ae2a146829f9 Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Sun, 17 Dec 2023 12:28:30 +0100 Subject: [PATCH 3/3] Update util/math_reference.h --- util/math_reference.h | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/util/math_reference.h b/util/math_reference.h index 533025aef..a2d455a93 100644 --- a/util/math_reference.h +++ b/util/math_reference.h @@ -309,10 +309,11 @@ sycl_cts::resultRef abs_diff(T a, T b) { using U = std::make_unsigned_t; T h = (a > b) ? a : b; T l = (a <= b) ? a : b; - // Use knowledge of 2-complement encoding while we know the difference is - // positive, so no overflow/underflow. + // Using two's-complement and that unsigned integer underflow is defined as + // modulo 2^n we get the result by computing the distance based on signed + // comparison. U result = static_cast(h) - static_cast(l); - return result > static_cast(std::numeric_limits::max()) + return result > std::numeric_limits::max() ? sycl_cts::resultRef(0, true) : T(result); }