[SYCL] Fix min/max known_identity<> wrong when -ffast-math is used. (#17028)

Technically a workaround, this fixes the problem with
`known_identity<minOp/maxOp>` being incorrect when using `-ffast-math`.
It works because we have a `__FAST_MATH__` macro that is defined when
that flag is used. Better, would be to detect `-fno-honor-infinities`.
Opening a JIRA/discussion for that, but this is a big improvement in the
interim.

See discussion #13813

Author: cperkinsintel

sycl/include/sycl/known_identity.hpp

@@ -258,10 +258,24 @@ 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)
+  static constexpr AccumulatorT value =
+      (std::numeric_limits<AccumulatorT>::max)();
+#else
   static constexpr AccumulatorT value = static_cast<AccumulatorT>(
       std::numeric_limits<AccumulatorT>::has_infinity
           ? std::numeric_limits<AccumulatorT>::infinity()
           : (std::numeric_limits<AccumulatorT>::max)());
+#endif
 };
 
 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
@@ -295,11 +309,19 @@ 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__
+  static constexpr AccumulatorT value =
+      (std::numeric_limits<AccumulatorT>::lowest)();
+#else
   static constexpr AccumulatorT value = static_cast<AccumulatorT>(
       std::numeric_limits<AccumulatorT>::has_infinity
           ? static_cast<AccumulatorT>(
                 -std::numeric_limits<AccumulatorT>::infinity())
           : std::numeric_limits<AccumulatorT>::lowest());
+#endif
 };
 
 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)

sycl/test/basic_tests/reduction/infinite_identity.cpp

@@ -0,0 +1,62 @@
+/*
+    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>
+    over {4.0f, 1.0f, 3.0f, 2.0f}  will return 0.0 instead of 1.0.
+*/
+
+// RUN: %clangxx -fsycl -fsyntax-only %s
+// RUN: %clangxx -fsycl -fsyntax-only %s -ffast-math
+// RUN: %clangxx -fsycl -fsyntax-only %s -fno-fast-math
+
+#include <sycl/ext/oneapi/bfloat16.hpp>
+#include <sycl/sycl.hpp>
+
+#include <cassert>
+#include <limits>
+
+template <typename OperandT> void test_known_identity_min() {
+  constexpr OperandT identity =
+      sycl::known_identity_v<sycl::minimum<OperandT>, OperandT>;
+#ifdef __FAST_MATH__
+  constexpr OperandT expected = std::numeric_limits<OperandT>::max();
+#else
+  constexpr OperandT expected = std::numeric_limits<OperandT>::infinity();
+#endif
+
+  static_assert(identity == expected,
+                "Known identity for sycl::minimum<T> is incorrect");
+}
+
+template <typename OperandT> void test_known_identity_max() {
+  constexpr OperandT identity =
+      sycl::known_identity_v<sycl::maximum<OperandT>, OperandT>;
+#ifdef __FAST_MATH__
+  constexpr OperandT expected = std::numeric_limits<OperandT>::lowest();
+#else
+  constexpr OperandT expected =
+      -std::numeric_limits<OperandT>::infinity(); // negative infinity
+#endif
+
+  static_assert(identity == expected,
+                "Known identity for sycl::maximum<T> is incorrect");
+}
+
+int main() {
+  test_known_identity_min<float>();
+  test_known_identity_min<double>();
+  test_known_identity_min<sycl::half>();
+
+  test_known_identity_max<float>();
+  test_known_identity_max<double>();
+  test_known_identity_max<sycl::half>();
+
+  // bfloat16 seems to be missing constexpr == which is used above.
+  // commenting out until fixed.
+  // test_known_identity_min<sycl::ext::oneapi::bfloat16>();
+  // test_known_identity_max<sycl::ext::oneapi::bfloat16>();
+
+  return 0;
+}