From 5c15647615e2fe973659afe0e0cb06393ba4a2f7 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Tue, 29 Apr 2025 11:44:35 -0700
Subject: [PATCH 1/3] tighten known_identity to finite_math macro

---
 sycl/include/sycl/known_identity.hpp          | 21 +++++++------------
 .../reduction/infinite_identity.cpp           | 12 +++++------
 2 files changed, 13 insertions(+), 20 deletions(-)

diff --git a/sycl/include/sycl/known_identity.hpp b/sycl/include/sycl/known_identity.hpp
index 91b05fcbd9b9b..f9ce91d4901eb 100644
--- a/sycl/include/sycl/known_identity.hpp
+++ b/sycl/include/sycl/known_identity.hpp
@@ -258,16 +258,11 @@ template <typename BinaryOperation, typename AccumulatorT>
 struct known_identity_impl<BinaryOperation, AccumulatorT,
                            std::enable_if_t<IsMinimumIdentityOp<
                                AccumulatorT, BinaryOperation>::value>> {
-// TODO: detect -fno-honor-infinities instead of -ffast-math
-// See https://github.com/intel/llvm/issues/13813
-// This workaround is a vast improvement, but still falls short.
-// To correct it properly, we need to detect -fno-honor-infinities usage,
-// perhaps via a driver inserted macro.
-// See similar below in known_identity_impl<IsMaximumIdentityOp>
-#ifdef __FAST_MATH__
-  // -ffast-math implies -fno-honor-infinities,
-  // but neither affect ::has_infinity (which is correct behavior, if
-  // unexpected)
+
+#if defined(__FINITE_MATH_ONLY__) && (__FINITE_MATH_ONLY__ == 1)
+  // Finite math only (-ffast-math,  -fno-honor-infinities) improves
+  // performance, but does not affect ::has_infinity (which is correct behavior,
+  // if unexpected). Use ::max() instead of ::infinity().
   static constexpr AccumulatorT value =
       (std::numeric_limits<AccumulatorT>::max)();
 #else
@@ -309,10 +304,8 @@ template <typename BinaryOperation, typename AccumulatorT>
 struct known_identity_impl<BinaryOperation, AccumulatorT,
                            std::enable_if_t<IsMaximumIdentityOp<
                                AccumulatorT, BinaryOperation>::value>> {
-// TODO: detect -fno-honor-infinities instead of -ffast-math
-// See https://github.com/intel/llvm/issues/13813
-// and comments above in known_identity_impl<IsMinimumIdentityOp>
-#ifdef __FAST_MATH__
+// See comment above in known_identity_impl<IsMinimumIdentityOp>
+#if defined(__FINITE_MATH_ONLY__) && (__FINITE_MATH_ONLY__ == 1)
   static constexpr AccumulatorT value =
       (std::numeric_limits<AccumulatorT>::lowest)();
 #else
diff --git a/sycl/test/basic_tests/reduction/infinite_identity.cpp b/sycl/test/basic_tests/reduction/infinite_identity.cpp
index 92eb9e35d2d44..f48678b77d663 100644
--- a/sycl/test/basic_tests/reduction/infinite_identity.cpp
+++ b/sycl/test/basic_tests/reduction/infinite_identity.cpp
@@ -1,6 +1,4 @@
 /*
-    See https://github.com/intel/llvm/issues/13813
-
     -ffast-math implies -fno-honor-infinities. In this case, the known_identity
    for sycl::minimum<T> cannot be inifinity (which will be 0), but must instead
     be the max<T> .  Otherwise, reducing sycl::minimum<T>
@@ -10,6 +8,8 @@
 // RUN: %clangxx -fsycl -fsyntax-only %s
 // RUN: %clangxx -fsycl -fsyntax-only %s -ffast-math
 // RUN: %clangxx -fsycl -fsyntax-only %s -fno-fast-math
+// RUN: %clangxx -fsycl -fsyntax-only %s -fno-fast-math -fno-honor-infinities -fno-honor-nans
+// RUN: %clangxx -fsycl -fsyntax-only %s -ffast-math -fhonor-infinities -fhonor-nans
 
 #include <sycl/ext/oneapi/bfloat16.hpp>
 #include <sycl/sycl.hpp>
@@ -20,7 +20,7 @@
 template <typename OperandT> void test_known_identity_min() {
   constexpr OperandT identity =
       sycl::known_identity_v<sycl::minimum<OperandT>, OperandT>;
-#ifdef __FAST_MATH__
+#if defined(__FINITE_MATH_ONLY__) && (__FINITE_MATH_ONLY__ == 1)
   constexpr OperandT expected = std::numeric_limits<OperandT>::max();
 #else
   constexpr OperandT expected = std::numeric_limits<OperandT>::infinity();
@@ -33,11 +33,11 @@ template <typename OperandT> void test_known_identity_min() {
 template <typename OperandT> void test_known_identity_max() {
   constexpr OperandT identity =
       sycl::known_identity_v<sycl::maximum<OperandT>, OperandT>;
-#ifdef __FAST_MATH__
+#if defined(__FINITE_MATH_ONLY__) && (__FINITE_MATH_ONLY__ == 1)
   constexpr OperandT expected = std::numeric_limits<OperandT>::lowest();
 #else
-  constexpr OperandT expected =
-      -std::numeric_limits<OperandT>::infinity(); // negative infinity
+  // negative infinity
+  constexpr OperandT expected = -std::numeric_limits<OperandT>::infinity();
 #endif
 
   static_assert(identity == expected,

From f849c41e7defb5674b96c518d78fc6e58eeb75e1 Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Wed, 30 Apr 2025 12:11:58 -0700
Subject: [PATCH 2/3] expanded comment

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 sycl/include/sycl/known_identity.hpp | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/sycl/include/sycl/known_identity.hpp b/sycl/include/sycl/known_identity.hpp
index f9ce91d4901eb..4dc48a7e88b80 100644
--- a/sycl/include/sycl/known_identity.hpp
+++ b/sycl/include/sycl/known_identity.hpp
@@ -263,6 +263,8 @@ struct known_identity_impl<BinaryOperation, AccumulatorT,
   // Finite math only (-ffast-math,  -fno-honor-infinities) improves
   // performance, but does not affect ::has_infinity (which is correct behavior,
   // if unexpected). Use ::max() instead of ::infinity().
+  // Note that if someone uses -fno-honor-infinities WITHOUT -fno-honor-nans
+  // they'll end up using ::infinity(). This is a very unlikely occurence
   static constexpr AccumulatorT value =
       (std::numeric_limits<AccumulatorT>::max)();
 #else

From 696797f9527dc649f9aca6ec1fa2e3d4fc200dcd Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Thu, 1 May 2025 13:25:37 -0700
Subject: [PATCH 3/3] care and feeding

Signed-off-by: Chris Perkins <chris.perkins@intel.com>
---
 sycl/include/sycl/known_identity.hpp | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/sycl/include/sycl/known_identity.hpp b/sycl/include/sycl/known_identity.hpp
index 4dc48a7e88b80..49a4305508c3f 100644
--- a/sycl/include/sycl/known_identity.hpp
+++ b/sycl/include/sycl/known_identity.hpp
@@ -264,7 +264,10 @@ struct known_identity_impl<BinaryOperation, AccumulatorT,
   // performance, but does not affect ::has_infinity (which is correct behavior,
   // if unexpected). Use ::max() instead of ::infinity().
   // Note that if someone uses -fno-honor-infinities WITHOUT -fno-honor-nans
-  // they'll end up using ::infinity(). This is a very unlikely occurence
+  // they'll end up using ::infinity(). There is no reasonable case where one of
+  // the flags would be used without the other and therefore
+  // __FINITE_MATH_ONLY__ is sufficient for this problem ( and superior to the
+  // __FAST_MATH__ macro we were using before).
   static constexpr AccumulatorT value =
       (std::numeric_limits<AccumulatorT>::max)();
 #else