[SYCL] tighten known_identity to finite_math macro (#18247)

Some time ago we fixed a bug in `sycl::known_identity` where it was
incorrectly using `std::numeric_limits<T>::infinity()` even in cases
where `::infinity()` should not have been available. At the time, the
fix was to simply test for `__FAST_MATH__`.

That worked, but was not ideal. Now we are tightening this to use the
`__FINITE_MATH_ONLY__` macro which is more accurate.
see #13813

`__FINITE_MATH_ONLY__` will be set when both the `-fno-honor-infinities`
and `-fno-honor-nans` flags are used. It was felt by the FE team that
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).

---------

Signed-off-by: Chris Perkins <chris.perkins@intel.com>

Author: cperkinsintel

sycl/include/sycl/known_identity.hpp

@@ -258,16 +258,16 @@ 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().
+  // Note that if someone uses -fno-honor-infinities WITHOUT -fno-honor-nans
+  // 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
@@ -309,10 +309,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

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,