diff --git a/sycl/doc/extensions/Reduction/Reduction.md b/sycl/doc/extensions/Reduction/Reduction.md index 37a6748142519..a454b2490faf0 100644 --- a/sycl/doc/extensions/Reduction/Reduction.md +++ b/sycl/doc/extensions/Reduction/Reduction.md @@ -38,6 +38,31 @@ unspecified reduction(span 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 +struct known_identity { + static constexpr AccumulatorT value; +}; + +// Available if C++17 +template +inline constexpr AccumulatorT known_identity_v = known_identity::value; +``` + +Whether `known_identity::value` exists can be tested using the `has_known_identity` trait class: + +```c++ +template +struct has_known_identity { + static constexpr bool value; +}; + +// Available if C++17 +template +inline constexpr bool has_known_identity_v = has_known_identity::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` and `std::atomic_ref`. diff --git a/sycl/include/CL/sycl/ONEAPI/reduction.hpp b/sycl/include/CL/sycl/ONEAPI/reduction.hpp index f15a811c21873..73ae2d7e5428b 100644 --- a/sycl/include/CL/sycl/ONEAPI/reduction.hpp +++ b/sycl/include/CL/sycl/ONEAPI/reduction.hpp @@ -144,6 +144,85 @@ using IsKnownIdentityOp = IsMinimumIdentityOp::value || IsMaximumIdentityOp::value>; +template +struct has_known_identity_impl + : std::integral_constant< + bool, IsKnownIdentityOp::value> {}; + +template +struct known_identity_impl {}; + +/// Returns zero as identity for ADD, OR, XOR operations. +template +struct known_identity_impl::value>::type> { + static constexpr AccumulatorT value = 0; +}; + +template +struct known_identity_impl::value>::type> { + static constexpr half value = +#ifdef __SYCL_DEVICE_ONLY__ + 0; +#else + cl::sycl::detail::host_half_impl::half(static_cast(0)); +#endif +}; + +/// Returns one as identify for MULTIPLY operations. +template +struct known_identity_impl::value>::type> { + static constexpr AccumulatorT value = 1; +}; + +template +struct known_identity_impl::value>::type> { + static constexpr half value = +#ifdef __SYCL_DEVICE_ONLY__ + 1; +#else + cl::sycl::detail::host_half_impl::half(static_cast(0x3C00)); +#endif +}; + +/// Returns bit image consisting of all ones as identity for AND operations. +template +struct known_identity_impl::value>::type> { + static constexpr AccumulatorT value = ~static_cast(0); +}; + +/// Returns maximal possible value as identity for MIN operations. +template +struct known_identity_impl::value>::type> { + static constexpr AccumulatorT value = + std::numeric_limits::has_infinity + ? std::numeric_limits::infinity() + : (std::numeric_limits::max)(); +}; + +/// Returns minimal possible value as identity for MAX operations. +template +struct known_identity_impl::value>::type> { + static constexpr AccumulatorT value = + std::numeric_limits::has_infinity + ? static_cast( + -std::numeric_limits::infinity()) + : std::numeric_limits::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 - static enable_if_t::value, _T> - getIdentity() { - return 0; - } - - /// Returns one as identify for MULTIPLY operations. - template - static enable_if_t::value, _T> - getIdentity() { - return 1; - } - - /// Returns bit image consisting of all ones as identity for AND operations. template - static enable_if_t::value, _T> + static enable_if_t::value, _T> getIdentity() { - return ~static_cast<_T>(0); - } - - /// Returns maximal possible value as identity for MIN operations. - template - static enable_if_t::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 - static enable_if_t::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 @@ -1062,6 +1108,26 @@ reduction(T *VarPtr, BinaryOperation) { access::mode::read_write>(VarPtr); } +template +struct has_known_identity : detail::has_known_identity_impl< + typename std::decay::type, + typename std::decay::type> {}; +#if __cplusplus >= 201703L +template +inline constexpr bool has_known_identity_v = + has_known_identity::value; +#endif + +template +struct known_identity + : detail::known_identity_impl::type, + typename std::decay::type> {}; +#if __cplusplus >= 201703L +template +inline constexpr AccumulatorT known_identity_v = + known_identity::value; +#endif + } // namespace ONEAPI } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/half_type.hpp b/sycl/include/CL/sycl/half_type.hpp index d7a0b22e09ee0..6627612f79cb2 100644 --- a/sycl/include/CL/sycl/half_type.hpp +++ b/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 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(Data); } @@ -280,8 +301,13 @@ template <> struct numeric_limits { 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(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 diff --git a/sycl/test/regression/constexpr-fp16-numeric-limits.cpp b/sycl/test/regression/constexpr-fp16-numeric-limits.cpp index feab488478683..c1e00be21b8bb 100644 --- a/sycl/test/regression/constexpr-fp16-numeric-limits.cpp +++ b/sycl/test/regression/constexpr-fp16-numeric-limits.cpp @@ -10,6 +10,8 @@ int main() { constexpr cl::sycl::half L5 = std::numeric_limits::round_error(); constexpr cl::sycl::half L6 = std::numeric_limits::infinity(); + constexpr cl::sycl::half L6n = + -std::numeric_limits::infinity(); constexpr cl::sycl::half L7 = std::numeric_limits::quiet_NaN(); constexpr cl::sycl::half L8 =