[SYCL] Add has_known_identity/known_identity (#2528)

* [SYCL] Add has_known_identity/known_identity type traits

These traits have been requested by developers who would like to
determine at compile-time whether the reduction() interface supports
their use-case, or whether they need to fall back to some other
implementation.

Signed-off-by: John Pennycook <john.pennycook@intel.com>
Co-authored-by: Roland Schulz <roland.schulz@intel.com>

Author: Pennycook

sycl/doc/extensions/Reduction/Reduction.md

@@ -38,6 +38,31 @@ unspecified reduction(span<T, Extent> var, const T& identity, BinaryOperation co
 
 The exact behavior of a reduction is specific to an implementation; the only interface exposed to the user is the set of functions above, which construct an unspecified `reduction` object encapsulating the reduction variable, an optional operator identity and the reduction operator.  For user-defined binary operations, an implementation should issue a compile-time warning if an identity is not specified and this is known to negatively impact performance (e.g. as a result of the implementation choosing a different reduction algorithm).  For standard binary operations (e.g. `std::plus`) on arithmetic types, the implementation must determine the correct identity automatically in order to avoid performance penalties.
 
+If an implementation can identify the identity value for a given combination of accumulator type `AccumulatorT` and function object type `BinaryOperation`, the value is defined as a member of the `known_identity` trait class:
+```c++
+template <typename BinaryOperation, typename AccumulatorT>
+struct known_identity {
+  static constexpr AccumulatorT value;
+};
+
+// Available if C++17
+template <typename BinaryOperation, typename AccumulatorT>
+inline constexpr AccumulatorT known_identity_v = known_identity<BinaryOperation, AccumulatorT>::value;
+```
+
+Whether `known_identity<BinaryOperation, AccumulatorT>::value` exists can be tested using the `has_known_identity` trait class:
+
+```c++
+template <typename BinaryOperation, typename AccumulatorT>
+struct has_known_identity {
+  static constexpr bool value;
+};
+
+// Available if C++17
+template <typename BinaryOperation, typename AccumulatorT>
+inline constexpr bool has_known_identity_v = has_known_identity<BinaryOperation, AccumulatorT>::value;
+```
+
 The dimensionality of the `accessor` passed to the `reduction` function specifies the dimensionality of the reduction variable: a 0-dimensional `accessor` represents a scalar reduction, and any other dimensionality represents an array reduction.  Specifying an array reduction of size N is functionally equivalent to specifying N independent scalar reductions.  The access mode of the accessor determines whether the reduction variable's original value is included in the reduction (i.e. for `access::mode::read_write` it is included, and for `access::mode::discard_write` it is not).  Multiple reductions aliasing the same output results in undefined behavior.
 
 `T` must be trivially copyable, permitting an implementation to (optionally) use atomic operations to implement the reduction.  This restriction is aligned with `std::atomic<T>` and `std::atomic_ref<T>`.

sycl/include/CL/sycl/ONEAPI/reduction.hpp

@@ -144,6 +144,85 @@ using IsKnownIdentityOp =
                   IsMinimumIdentityOp<T, BinaryOperation>::value ||
                   IsMaximumIdentityOp<T, BinaryOperation>::value>;
 
+template <typename BinaryOperation, typename AccumulatorT>
+struct has_known_identity_impl
+    : std::integral_constant<
+          bool, IsKnownIdentityOp<AccumulatorT, BinaryOperation>::value> {};
+
+template <typename BinaryOperation, typename AccumulatorT, typename = void>
+struct known_identity_impl {};
+
+/// Returns zero as identity for ADD, OR, XOR operations.
+template <typename BinaryOperation, typename AccumulatorT>
+struct known_identity_impl<BinaryOperation, AccumulatorT,
+                           typename std::enable_if<IsZeroIdentityOp<
+                               AccumulatorT, BinaryOperation>::value>::type> {
+  static constexpr AccumulatorT value = 0;
+};
+
+template <typename BinaryOperation>
+struct known_identity_impl<BinaryOperation, half,
+                           typename std::enable_if<IsZeroIdentityOp<
+                               half, BinaryOperation>::value>::type> {
+  static constexpr half value =
+#ifdef __SYCL_DEVICE_ONLY__
+      0;
+#else
+      cl::sycl::detail::host_half_impl::half(static_cast<uint16_t>(0));
+#endif
+};
+
+/// Returns one as identify for MULTIPLY operations.
+template <typename BinaryOperation, typename AccumulatorT>
+struct known_identity_impl<BinaryOperation, AccumulatorT,
+                           typename std::enable_if<IsOneIdentityOp<
+                               AccumulatorT, BinaryOperation>::value>::type> {
+  static constexpr AccumulatorT value = 1;
+};
+
+template <typename BinaryOperation>
+struct known_identity_impl<BinaryOperation, half,
+                           typename std::enable_if<IsOneIdentityOp<
+                               half, BinaryOperation>::value>::type> {
+  static constexpr half value =
+#ifdef __SYCL_DEVICE_ONLY__
+      1;
+#else
+      cl::sycl::detail::host_half_impl::half(static_cast<uint16_t>(0x3C00));
+#endif
+};
+
+/// Returns bit image consisting of all ones as identity for AND operations.
+template <typename BinaryOperation, typename AccumulatorT>
+struct known_identity_impl<BinaryOperation, AccumulatorT,
+                           typename std::enable_if<IsOnesIdentityOp<
+                               AccumulatorT, BinaryOperation>::value>::type> {
+  static constexpr AccumulatorT value = ~static_cast<AccumulatorT>(0);
+};
+
+/// Returns maximal possible value as identity for MIN operations.
+template <typename BinaryOperation, typename AccumulatorT>
+struct known_identity_impl<BinaryOperation, AccumulatorT,
+                           typename std::enable_if<IsMinimumIdentityOp<
+                               AccumulatorT, BinaryOperation>::value>::type> {
+  static constexpr AccumulatorT value =
+      std::numeric_limits<AccumulatorT>::has_infinity
+          ? std::numeric_limits<AccumulatorT>::infinity()
+          : (std::numeric_limits<AccumulatorT>::max)();
+};
+
+/// Returns minimal possible value as identity for MAX operations.
+template <typename BinaryOperation, typename AccumulatorT>
+struct known_identity_impl<BinaryOperation, AccumulatorT,
+                           typename std::enable_if<IsMaximumIdentityOp<
+                               AccumulatorT, BinaryOperation>::value>::type> {
+  static constexpr AccumulatorT value =
+      std::numeric_limits<AccumulatorT>::has_infinity
+          ? static_cast<AccumulatorT>(
+                -std::numeric_limits<AccumulatorT>::infinity())
+          : std::numeric_limits<AccumulatorT>::lowest();
+};
+
 /// Class that is used to represent objects that are passed to user's lambda
 /// functions and representing users' reduction variable.
 /// The generic version of the class represents those reductions of those
@@ -193,43 +272,10 @@ class reducer<T, BinaryOperation,
     MValue = BOp(MValue, Partial);
   }
 
-  /// Returns zero as identity for ADD, OR, XOR operations.
-  template <typename _T = T, class _BinaryOperation = BinaryOperation>
-  static enable_if_t<IsZeroIdentityOp<_T, _BinaryOperation>::value, _T>
-  getIdentity() {
-    return 0;
-  }
-
-  /// Returns one as identify for MULTIPLY operations.
-  template <typename _T = T, class _BinaryOperation = BinaryOperation>
-  static enable_if_t<IsOneIdentityOp<_T, _BinaryOperation>::value, _T>
-  getIdentity() {
-    return 1;
-  }
-
-  /// Returns bit image consisting of all ones as identity for AND operations.
   template <typename _T = T, class _BinaryOperation = BinaryOperation>
-  static enable_if_t<IsOnesIdentityOp<_T, _BinaryOperation>::value, _T>
+  static enable_if_t<has_known_identity_impl<_BinaryOperation, _T>::value, _T>
   getIdentity() {
-    return ~static_cast<_T>(0);
-  }
-
-  /// Returns maximal possible value as identity for MIN operations.
-  template <typename _T = T, class _BinaryOperation = BinaryOperation>
-  static enable_if_t<IsMinimumIdentityOp<_T, _BinaryOperation>::value, _T>
-  getIdentity() {
-    return std::numeric_limits<_T>::has_infinity
-               ? std::numeric_limits<_T>::infinity()
-               : (std::numeric_limits<_T>::max)();
-  }
-
-  /// Returns minimal possible value as identity for MAX operations.
-  template <typename _T = T, class _BinaryOperation = BinaryOperation>
-  static enable_if_t<IsMaximumIdentityOp<_T, _BinaryOperation>::value, _T>
-  getIdentity() {
-    return std::numeric_limits<_T>::has_infinity
-               ? static_cast<_T>(-std::numeric_limits<_T>::infinity())
-               : std::numeric_limits<_T>::lowest();
+    return known_identity_impl<_BinaryOperation, _T>::value;
   }
 
   template <typename _T = T>
@@ -1062,6 +1108,26 @@ reduction(T *VarPtr, BinaryOperation) {
                                 access::mode::read_write>(VarPtr);
 }
 
+template <typename BinaryOperation, typename AccumulatorT>
+struct has_known_identity : detail::has_known_identity_impl<
+                                typename std::decay<BinaryOperation>::type,
+                                typename std::decay<AccumulatorT>::type> {};
+#if __cplusplus >= 201703L
+template <typename BinaryOperation, typename AccumulatorT>
+inline constexpr bool has_known_identity_v =
+    has_known_identity<BinaryOperation, AccumulatorT>::value;
+#endif
+
+template <typename BinaryOperation, typename AccumulatorT>
+struct known_identity
+    : detail::known_identity_impl<typename std::decay<BinaryOperation>::type,
+                                  typename std::decay<AccumulatorT>::type> {};
+#if __cplusplus >= 201703L
+template <typename BinaryOperation, typename AccumulatorT>
+inline constexpr AccumulatorT known_identity_v =
+    known_identity<BinaryOperation, AccumulatorT>::value;
+#endif
+
 } // namespace ONEAPI
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/half_type.hpp

@@ -26,6 +26,11 @@
 #else
 #define __SYCL_CONSTEXPR_ON_DEVICE
 #endif
+#if __cplusplus >= 201402L
+#define _CPP14_CONSTEXPR constexpr
+#else
+#define _CPP14_CONSTEXPR
+#endif
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
@@ -35,8 +40,8 @@ namespace host_half_impl {
 class __SYCL_EXPORT half {
 public:
   half() = default;
-  half(const half &) = default;
-  half(half &&) = default;
+  constexpr half(const half &) = default;
+  constexpr half(half &&) = default;
 
   half(const float &rhs);
 
@@ -74,11 +79,20 @@ class __SYCL_EXPORT half {
     return ret;
   }
 
+  // Operator neg
+  _CPP14_CONSTEXPR half &operator-() {
+    Buf ^= 0x8000;
+    return *this;
+  }
+
   // Operator float
   operator float() const;
 
   template <typename Key> friend struct std::hash;
 
+  // Initialize underlying data
+  constexpr explicit half(uint16_t x) : Buf(x) {}
+
 private:
   uint16_t Buf;
 };
@@ -136,8 +150,8 @@ class half;
 class half {
 public:
   half() = default;
-  half(const half &) = default;
-  half(half &&) = default;
+  constexpr half(const half &) = default;
+  constexpr half(half &&) = default;
 
   __SYCL_CONSTEXPR_ON_DEVICE half(const float &rhs) : Data(rhs) {}
 
@@ -146,8 +160,8 @@ class half {
 #ifndef __SYCL_DEVICE_ONLY__
   // Since StorageT and BIsRepresentationT are different on host, these two
   // helpers are required for 'vec' class
-  half(const detail::host_half_impl::half &rhs) : Data(rhs) {};
-  operator detail::host_half_impl::half() const { return Data; }
+  constexpr half(const detail::host_half_impl::half &rhs) : Data(rhs){};
+  constexpr operator detail::host_half_impl::half() const { return Data; }
 #endif // __SYCL_DEVICE_ONLY__
 
   // Operator +=, -=, *=, /=
@@ -193,7 +207,14 @@ class half {
     operator--();
     return ret;
   }
-
+  _CPP14_CONSTEXPR half &operator-() {
+    Data = -Data;
+    return *this;
+  }
+  _CPP14_CONSTEXPR half operator-() const {
+    half r = *this;
+    return -r;
+  }
   // Operator float
   operator float() const { return static_cast<float>(Data); }
 
@@ -280,8 +301,13 @@ template <> struct numeric_limits<cl::sycl::half> {
     return 0.5f;
   }
 
-  static __SYCL_CONSTEXPR_ON_DEVICE const cl::sycl::half infinity() noexcept {
+  static constexpr const cl::sycl::half infinity() noexcept {
+#ifdef __SYCL_DEVICE_ONLY__
     return __builtin_huge_valf();
+#else
+    return cl::sycl::detail::host_half_impl::half(
+        static_cast<uint16_t>(0x7C00));
+#endif
   }
 
   static __SYCL_CONSTEXPR_ON_DEVICE const cl::sycl::half quiet_NaN() noexcept {
@@ -313,3 +339,4 @@ inline std::istream &operator>>(std::istream &I, cl::sycl::half &rhs) {
 }
 
 #undef __SYCL_CONSTEXPR_ON_DEVICE
+#undef _CPP14_CONSTEXPR

sycl/test/regression/constexpr-fp16-numeric-limits.cpp

@@ -10,6 +10,8 @@ int main() {
   constexpr cl::sycl::half L5 =
       std::numeric_limits<cl::sycl::half>::round_error();
   constexpr cl::sycl::half L6 = std::numeric_limits<cl::sycl::half>::infinity();
+  constexpr cl::sycl::half L6n =
+      -std::numeric_limits<cl::sycl::half>::infinity();
   constexpr cl::sycl::half L7 =
       std::numeric_limits<cl::sycl::half>::quiet_NaN();
   constexpr cl::sycl::half L8 =