From 9949d126f9443b725fcba4f4666b8c30f2fa5831 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 13 Nov 2025 08:23:22 -0800 Subject: [PATCH 1/6] [SYCL] Fix complex tanh The current implementation of the tanh function for the SYCL complex extension did not always return values precise enough, neither on device nor on host. This was caused by accumulated error from a call to cos. This commit changes the implementation of tanh to use a different way of calculating the result and tests it against std::tanh for complex numbers. Signed-off-by: Larsen, Steffen --- .../complex/detail/complex_math.hpp | 24 ++--- sycl/test/regression/tanh_complex.cpp | 100 ++++++++++++++++++ 2 files changed, 111 insertions(+), 13 deletions(-) create mode 100644 sycl/test/regression/tanh_complex.cpp diff --git a/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex_math.hpp b/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex_math.hpp index 0d84fc5060a3f..8e3616d9feb0e 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex_math.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex_math.hpp @@ -493,22 +493,20 @@ template __DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY typename std::enable_if_t::value, complex<_Tp>> tanh(const complex<_Tp> &__x) { - if (sycl::isinf(__x.real())) { - if (!sycl::isfinite(__x.imag())) - return complex<_Tp>(sycl::copysign(_Tp(1), __x.real()), _Tp(0)); + if (sycl::isinf(__x.real())) return complex<_Tp>(sycl::copysign(_Tp(1), __x.real()), - sycl::copysign(_Tp(0), sycl::sin(_Tp(2) * __x.imag()))); - } + sycl::copysign(_Tp(0), sycl::isfinite(__x.imag()) + ? sin(_Tp(2) * __x.imag()) + : _Tp(1))); if (sycl::isnan(__x.real()) && __x.imag() == 0) return __x; - _Tp __2r(_Tp(2) * __x.real()); - _Tp __2i(_Tp(2) * __x.imag()); - _Tp __d(sycl::cosh(__2r) + sycl::cos(__2i)); - _Tp __2rsh(sycl::sinh(__2r)); - if (sycl::isinf(__2rsh) && sycl::isinf(__d)) - return complex<_Tp>(__2rsh > _Tp(0) ? _Tp(1) : _Tp(-1), - __2i > _Tp(0) ? _Tp(0) : _Tp(-0.)); - return complex<_Tp>(__2rsh / __d, sycl::sin(__2i) / __d); + complex<_Tp> sinh_x = sinh(__x); + complex<_Tp> cosh_x = cosh(__x); + if (sycl::isinf(sinh_x.real()) && sycl::isinf(cosh_x.real())) + return complex<_Tp>(sinh_x.real() * cosh_x.real() > _Tp(0) ? _Tp(1) + : _Tp(-1), + __x.imag() > _Tp(0) ? _Tp(0) : _Tp(-0.)); + return sinh_x / cosh_x; } // asin diff --git a/sycl/test/regression/tanh_complex.cpp b/sycl/test/regression/tanh_complex.cpp new file mode 100644 index 0000000000000..444b95a2201bf --- /dev/null +++ b/sycl/test/regression/tanh_complex.cpp @@ -0,0 +1,100 @@ +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: %t.out + +// Checks the results of tanh on certain complex numbers. + +#define SYCL_EXT_ONEAPI_COMPLEX + +#include +#include + +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +int Failures = 0; + +template bool FloatingPointEq(T LHS, T RHS) { + if (std::isnan(LHS)) + return std::isnan(RHS); + return LHS == RHS; +} + +#define CHECK_TANH_RESULT(REAL, IMAG, T) \ + { \ + syclexp::complex sycl_res = \ + syclexp::tanh(syclexp::complex{REAL, IMAG}); \ + std::complex std_res = std::tanh(std::complex{REAL, IMAG}); \ + if (!FloatingPointEq(sycl_res.real(), std_res.real())) { \ + std::cout << "Real differ in tanh((" << REAL << ", " << IMAG \ + << ")): " << sycl_res.real() << " != " << std_res.real() \ + << std::endl; \ + ++Failures; \ + } \ + if (!FloatingPointEq(sycl_res.imag(), std_res.imag())) { \ + std::cout << "Imag differ in tanh((" << REAL << ", " << IMAG \ + << ")): " << sycl_res.imag() << " != " << std_res.imag() \ + << std::endl; \ + ++Failures; \ + } \ + } + +int main() { + CHECK_TANH_RESULT(0, -11.0, float); + CHECK_TANH_RESULT(0, -11.0, double); + + CHECK_TANH_RESULT(std::numeric_limits::infinity(), 32.0, float); + CHECK_TANH_RESULT(std::numeric_limits::infinity(), 32.0, double); + + CHECK_TANH_RESULT(32.0, std::numeric_limits::infinity(), float); + CHECK_TANH_RESULT(32.0, std::numeric_limits::infinity(), double); + + CHECK_TANH_RESULT(std::numeric_limits::infinity(), + std::numeric_limits::infinity(), float); + CHECK_TANH_RESULT(std::numeric_limits::infinity(), + std::numeric_limits::infinity(), double); + + CHECK_TANH_RESULT(-std::numeric_limits::infinity(), 32.0, float); + CHECK_TANH_RESULT(-std::numeric_limits::infinity(), 32.0, double); + + CHECK_TANH_RESULT(32.0, -std::numeric_limits::infinity(), float); + CHECK_TANH_RESULT(32.0, -std::numeric_limits::infinity(), double); + + CHECK_TANH_RESULT(-std::numeric_limits::infinity(), + -std::numeric_limits::infinity(), float); + CHECK_TANH_RESULT(-std::numeric_limits::infinity(), + -std::numeric_limits::infinity(), double); + + CHECK_TANH_RESULT(std::numeric_limits::max(), 0.0, float); + CHECK_TANH_RESULT(std::numeric_limits::max(), 0.0, double); + + CHECK_TANH_RESULT(0.0, std::numeric_limits::max(), float); + CHECK_TANH_RESULT(0.0, std::numeric_limits::max(), double); + + CHECK_TANH_RESULT(0.0, 0.0, float); + CHECK_TANH_RESULT(0.0, 0.0, double); + + CHECK_TANH_RESULT(std::numeric_limits::infinity(), + std::numeric_limits::quiet_NaN(), float); + CHECK_TANH_RESULT(std::numeric_limits::infinity(), + std::numeric_limits::quiet_NaN(), double); + + CHECK_TANH_RESULT(-std::numeric_limits::infinity(), + std::numeric_limits::quiet_NaN(), float); + CHECK_TANH_RESULT(-std::numeric_limits::infinity(), + std::numeric_limits::quiet_NaN(), double); + + CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), 0.0, float); + CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), 0.0, double); + + CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), 1.0, float); + CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), 1.0, double); + + CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), + std::numeric_limits::quiet_NaN(), float); + CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), + std::numeric_limits::quiet_NaN(), double); + + return Failures; +} From 4b56ea6e296af31003a78fa9a147c9c122e76918 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 14 Nov 2025 00:51:30 -0800 Subject: [PATCH 2/6] Fix namespace qualification Signed-off-by: Larsen, Steffen --- .../oneapi/experimental/complex/detail/complex_math.hpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex_math.hpp b/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex_math.hpp index 8e3616d9feb0e..c9557f83d2093 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex_math.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex_math.hpp @@ -494,10 +494,11 @@ __DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY typename std::enable_if_t::value, complex<_Tp>> tanh(const complex<_Tp> &__x) { if (sycl::isinf(__x.real())) - return complex<_Tp>(sycl::copysign(_Tp(1), __x.real()), - sycl::copysign(_Tp(0), sycl::isfinite(__x.imag()) - ? sin(_Tp(2) * __x.imag()) - : _Tp(1))); + return complex<_Tp>( + sycl::copysign(_Tp(1), __x.real()), + sycl::copysign(_Tp(0), sycl::isfinite(__x.imag()) + ? sycl::sin(_Tp(2) * __x.imag()) + : _Tp(1))); if (sycl::isnan(__x.real()) && __x.imag() == 0) return __x; complex<_Tp> sinh_x = sinh(__x); From 218b25ed39fe8e9de433b1231237451a99c3a34a Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 14 Nov 2025 02:57:17 -0800 Subject: [PATCH 3/6] Address Windows bug and be lenient on precision Signed-off-by: Larsen, Steffen --- sycl/test/regression/tanh_complex.cpp | 88 +++++++++++++++++++-------- 1 file changed, 61 insertions(+), 27 deletions(-) diff --git a/sycl/test/regression/tanh_complex.cpp b/sycl/test/regression/tanh_complex.cpp index 444b95a2201bf..1c7b030d02e31 100644 --- a/sycl/test/regression/tanh_complex.cpp +++ b/sycl/test/regression/tanh_complex.cpp @@ -18,7 +18,8 @@ int Failures = 0; template bool FloatingPointEq(T LHS, T RHS) { if (std::isnan(LHS)) return std::isnan(RHS); - return LHS == RHS; + // Allow some rounding differences, but minimal. + return std::abs(LHS - RHS) < T{0.000001}; } #define CHECK_TANH_RESULT(REAL, IMAG, T) \ @@ -40,32 +41,34 @@ template bool FloatingPointEq(T LHS, T RHS) { } \ } +#define CHECK_TANH_REF_RESULT(REAL, IMAG, REF_REAL, REF_IMAG, T) \ + { \ + syclexp::complex sycl_res = \ + syclexp::tanh(syclexp::complex{REAL, IMAG}); \ + if (!FloatingPointEq(sycl_res.real(), T{REF_REAL})) { \ + std::cout << "Real differ in tanh((" << REAL << ", " << IMAG \ + << ")): " << sycl_res.real() << " != " << REF_REAL \ + << std::endl; \ + ++Failures; \ + } \ + if (!FloatingPointEq(sycl_res.imag(), T{REF_IMAG})) { \ + std::cout << "Imag differ in tanh((" << REAL << ", " << IMAG \ + << ")): " << sycl_res.imag() << " != " << REF_IMAG \ + << std::endl; \ + ++Failures; \ + } \ + } + int main() { CHECK_TANH_RESULT(0, -11.0, float); CHECK_TANH_RESULT(0, -11.0, double); - CHECK_TANH_RESULT(std::numeric_limits::infinity(), 32.0, float); - CHECK_TANH_RESULT(std::numeric_limits::infinity(), 32.0, double); - CHECK_TANH_RESULT(32.0, std::numeric_limits::infinity(), float); CHECK_TANH_RESULT(32.0, std::numeric_limits::infinity(), double); - CHECK_TANH_RESULT(std::numeric_limits::infinity(), - std::numeric_limits::infinity(), float); - CHECK_TANH_RESULT(std::numeric_limits::infinity(), - std::numeric_limits::infinity(), double); - - CHECK_TANH_RESULT(-std::numeric_limits::infinity(), 32.0, float); - CHECK_TANH_RESULT(-std::numeric_limits::infinity(), 32.0, double); - CHECK_TANH_RESULT(32.0, -std::numeric_limits::infinity(), float); CHECK_TANH_RESULT(32.0, -std::numeric_limits::infinity(), double); - CHECK_TANH_RESULT(-std::numeric_limits::infinity(), - -std::numeric_limits::infinity(), float); - CHECK_TANH_RESULT(-std::numeric_limits::infinity(), - -std::numeric_limits::infinity(), double); - CHECK_TANH_RESULT(std::numeric_limits::max(), 0.0, float); CHECK_TANH_RESULT(std::numeric_limits::max(), 0.0, double); @@ -75,16 +78,6 @@ int main() { CHECK_TANH_RESULT(0.0, 0.0, float); CHECK_TANH_RESULT(0.0, 0.0, double); - CHECK_TANH_RESULT(std::numeric_limits::infinity(), - std::numeric_limits::quiet_NaN(), float); - CHECK_TANH_RESULT(std::numeric_limits::infinity(), - std::numeric_limits::quiet_NaN(), double); - - CHECK_TANH_RESULT(-std::numeric_limits::infinity(), - std::numeric_limits::quiet_NaN(), float); - CHECK_TANH_RESULT(-std::numeric_limits::infinity(), - std::numeric_limits::quiet_NaN(), double); - CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), 0.0, float); CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), 0.0, double); @@ -96,5 +89,46 @@ int main() { CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), std::numeric_limits::quiet_NaN(), double); + // The MSVC implementation of tanh for complex numbers does not return + // (1, +-0) for complex numbers with an infinite real, despite the C++ + // standard requiring it. So instead, we check the results using reference + // values rather than trusting the result of std::tanh. + CHECK_TANH_REF_RESULT(std::numeric_limits::infinity(), 32.0, 1.0, 0.0, + float); + CHECK_TANH_REF_RESULT(std::numeric_limits::infinity(), 32.0, 1.0, 0.0, + double); + + CHECK_TANH_REF_RESULT(std::numeric_limits::infinity(), + std::numeric_limits::infinity(), 1, 0.0, float); + CHECK_TANH_REF_RESULT(std::numeric_limits::infinity(), + std::numeric_limits::infinity(), 1, 0.0, + double); + + CHECK_TANH_REF_RESULT(-std::numeric_limits::infinity(), 32.0, -1.0, + 0.0, float); + CHECK_TANH_REF_RESULT(-std::numeric_limits::infinity(), 32.0, -1.0, + 0.0, double); + + CHECK_TANH_REF_RESULT(-std::numeric_limits::infinity(), + -std::numeric_limits::infinity(), -1.0, 0.0, + float); + CHECK_TANH_REF_RESULT(-std::numeric_limits::infinity(), + -std::numeric_limits::infinity(), -1.0, 0.0, + double); + + CHECK_TANH_REF_RESULT(std::numeric_limits::infinity(), + std::numeric_limits::quiet_NaN(), 1.0, 0.0, + float); + CHECK_TANH_REF_RESULT(std::numeric_limits::infinity(), + std::numeric_limits::quiet_NaN(), 1.0, 0.0, + double); + + CHECK_TANH_REF_RESULT(-std::numeric_limits::infinity(), + std::numeric_limits::quiet_NaN(), -1.0, 0.0, + float); + CHECK_TANH_REF_RESULT(-std::numeric_limits::infinity(), + std::numeric_limits::quiet_NaN(), -1.0, 0.0, + double); + return Failures; } From 3f5e5965211f8b4d419899a6af43057e07734284 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 14 Nov 2025 03:49:15 -0800 Subject: [PATCH 4/6] Add another case for Windows references Signed-off-by: Larsen, Steffen --- sycl/test/regression/tanh_complex.cpp | 21 +++++++++++++-------- 1 file changed, 13 insertions(+), 8 deletions(-) diff --git a/sycl/test/regression/tanh_complex.cpp b/sycl/test/regression/tanh_complex.cpp index 1c7b030d02e31..be29d6cd742ee 100644 --- a/sycl/test/regression/tanh_complex.cpp +++ b/sycl/test/regression/tanh_complex.cpp @@ -19,7 +19,7 @@ template bool FloatingPointEq(T LHS, T RHS) { if (std::isnan(LHS)) return std::isnan(RHS); // Allow some rounding differences, but minimal. - return std::abs(LHS - RHS) < T{0.000001}; + return std::abs(LHS - RHS) < T{0.00001}; } #define CHECK_TANH_RESULT(REAL, IMAG, T) \ @@ -78,9 +78,6 @@ int main() { CHECK_TANH_RESULT(0.0, 0.0, float); CHECK_TANH_RESULT(0.0, 0.0, double); - CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), 0.0, float); - CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), 0.0, double); - CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), 1.0, float); CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), 1.0, double); @@ -89,10 +86,13 @@ int main() { CHECK_TANH_RESULT(std::numeric_limits::quiet_NaN(), std::numeric_limits::quiet_NaN(), double); - // The MSVC implementation of tanh for complex numbers does not return - // (1, +-0) for complex numbers with an infinite real, despite the C++ - // standard requiring it. So instead, we check the results using reference - // values rather than trusting the result of std::tanh. + // The MSVC implementation of tanh for complex numbers does not adhere to the + // following requirements set by the definition of std::tanh: + // * When the input has an infinite real, then the function should return + // (1, +-0). + // * When the input is (NaN, 0), the result should be (NaN, 0). + // Instead we check the results using reference values rather than trusting + // the result of std::tanh in these cases. CHECK_TANH_REF_RESULT(std::numeric_limits::infinity(), 32.0, 1.0, 0.0, float); CHECK_TANH_REF_RESULT(std::numeric_limits::infinity(), 32.0, 1.0, 0.0, @@ -130,5 +130,10 @@ int main() { std::numeric_limits::quiet_NaN(), -1.0, 0.0, double); + CHECK_TANH_REF_RESULT(std::numeric_limits::quiet_NaN(), 0.0, + std::numeric_limits::quiet_NaN(), 0.0, float); + CHECK_TANH_REF_RESULT(std::numeric_limits::quiet_NaN(), 0.0, + std::numeric_limits::quiet_NaN(), 0.0, double); + return Failures; } From 32557d6ac3a596ef0be77b09e682cd16d84d96a5 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 14 Nov 2025 04:49:17 -0800 Subject: [PATCH 5/6] Add precision for easier debugging Signed-off-by: Larsen, Steffen --- sycl/test/regression/tanh_complex.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/test/regression/tanh_complex.cpp b/sycl/test/regression/tanh_complex.cpp index be29d6cd742ee..9a183f0fe2edb 100644 --- a/sycl/test/regression/tanh_complex.cpp +++ b/sycl/test/regression/tanh_complex.cpp @@ -60,6 +60,9 @@ template bool FloatingPointEq(T LHS, T RHS) { } int main() { + // Set precision for easier debugging. + std::cout << std::setprecision(10); + CHECK_TANH_RESULT(0, -11.0, float); CHECK_TANH_RESULT(0, -11.0, double); From 5da51cb164936d30aa24413dec799098f818131f Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 14 Nov 2025 06:55:46 -0800 Subject: [PATCH 6/6] Relax precision as MSVC is off Signed-off-by: Larsen, Steffen --- sycl/test/regression/tanh_complex.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/regression/tanh_complex.cpp b/sycl/test/regression/tanh_complex.cpp index 9a183f0fe2edb..3396cf0da2004 100644 --- a/sycl/test/regression/tanh_complex.cpp +++ b/sycl/test/regression/tanh_complex.cpp @@ -19,7 +19,7 @@ template bool FloatingPointEq(T LHS, T RHS) { if (std::isnan(LHS)) return std::isnan(RHS); // Allow some rounding differences, but minimal. - return std::abs(LHS - RHS) < T{0.00001}; + return std::abs(LHS - RHS) < T{0.0001}; } #define CHECK_TANH_RESULT(REAL, IMAG, T) \