Skip to content

[SYCL] tighten known_identity to finite_math macro #18247

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 3 commits into
base: sycl
Choose a base branch
from
Open

[SYCL] tighten known_identity to finite_math macro #18247

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
5c15647
tighten known_identity to finite_math macro
cperkinsintel Apr 29, 2025
f849c41
expanded comment
cperkinsintel Apr 30, 2025
696797f
care and feeding
cperkinsintel May 1, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
26 changes: 12 additions & 14 deletions sycl/include/sycl/known_identity.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand Down
12 changes: 6 additions & 6 deletions sycl/test/basic_tests/reduction/infinite_identity.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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>
Expand All @@ -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>
Expand All @@ -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();
Expand All @@ -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,
Expand Down