diff --git a/sycl/include/sycl/ext/oneapi/experimental/complex/complex.hpp b/sycl/include/sycl/ext/oneapi/experimental/complex/complex.hpp new file mode 100644 index 0000000000000..ad6653081ff48 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/experimental/complex/complex.hpp @@ -0,0 +1,16 @@ +//===- complex.hpp --------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#ifdef SYCL_EXT_ONEAPI_COMPLEX + +#include "./detail/complex.hpp" +#include "./detail/complex_math.hpp" + +#endif // SYCL_EXT_ONEAPI_COMPLEX diff --git a/sycl/include/sycl/ext/oneapi/experimental/complex/detail/common.hpp b/sycl/include/sycl/ext/oneapi/experimental/complex/detail/common.hpp new file mode 100644 index 0000000000000..8e34cb5136982 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/experimental/complex/detail/common.hpp @@ -0,0 +1,62 @@ +//===- common.hpp ---------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include + +#include + +namespace sycl { +inline namespace _V1 { + +namespace ext { +namespace oneapi { +namespace experimental { + +//////////////////////////////////////////////////////////////////////////////// +/// FORWARD DECLARATIONS +//////////////////////////////////////////////////////////////////////////////// + +template struct is_genfloat; + +template class complex; + +template +class complex<_Tp, typename std::enable_if_t::value>>; + +//////////////////////////////////////////////////////////////////////////////// +/// TRAITS +//////////////////////////////////////////////////////////////////////////////// + +template +struct is_genfloat + : std::integral_constant || + std::is_same_v<_Tp, float> || + std::is_same_v<_Tp, sycl::half>> {}; + +template +struct is_gencomplex + : std::integral_constant> || + std::is_same_v<_Tp, complex> || + std::is_same_v<_Tp, complex>> {}; + +//////////////////////////////////////////////////////////////////////////////// +/// DEFINES +//////////////////////////////////////////////////////////////////////////////// + +#define _SYCL_EXT_CPLX_INLINE_VISIBILITY \ + inline __attribute__((__visibility__("hidden"), __always_inline__)) + +} // namespace experimental +} // namespace oneapi +} // namespace ext + +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex.hpp b/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex.hpp new file mode 100644 index 0000000000000..9557aa0e01b29 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex.hpp @@ -0,0 +1,388 @@ +//===- complex.hpp --------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include "common.hpp" + +#include + +#include + +namespace sycl { +inline namespace _V1 { + +namespace ext { +namespace oneapi { +namespace experimental { + +template +class complex<_Tp, typename std::enable_if_t::value>> { +public: + typedef _Tp value_type; + +private: + value_type __re_; + value_type __im_; + +public: + _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr complex( + value_type __re = value_type(), value_type __im = value_type()) + : __re_(__re), __im_(__im) {} + + template + _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr complex(const complex<_Xp> &__c) + : __re_(__c.real()), __im_(__c.imag()) {} + + template ::value>> + _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr complex( + const std::complex<_Xp> &__c) + : __re_(static_cast(__c.real())), + __im_(static_cast(__c.imag())) {} + + template ::value>> + _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr + operator std::complex<_Xp>() const { + return std::complex<_Xp>(static_cast<_Xp>(__re_), static_cast<_Xp>(__im_)); + } + + _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr value_type real() const { + return __re_; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr value_type imag() const { + return __im_; + } + + _SYCL_EXT_CPLX_INLINE_VISIBILITY void real(value_type __re) { __re_ = __re; } + _SYCL_EXT_CPLX_INLINE_VISIBILITY void imag(value_type __im) { __im_ = __im; } + + _SYCL_EXT_CPLX_INLINE_VISIBILITY complex &operator=(value_type __re) { + __re_ = __re; + __im_ = value_type(); + return *this; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & + operator+=(complex &__c, value_type __re) { + __c.__re_ += __re; + return __c; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & + operator-=(complex &__c, value_type __re) { + __c.__re_ -= __re; + return __c; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & + operator*=(complex &__c, value_type __re) { + __c.__re_ *= __re; + __c.__im_ *= __re; + return __c; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & + operator/=(complex &__c, value_type __re) { + __c.__re_ /= __re; + __c.__im_ /= __re; + return __c; + } + + template + _SYCL_EXT_CPLX_INLINE_VISIBILITY complex &operator=(const complex<_Xp> &__c) { + __re_ = __c.real(); + __im_ = __c.imag(); + return *this; + } + template + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & + operator+=(complex &__x, const complex<_Xp> &__y) { + __x.__re_ += __y.real(); + __x.__im_ += __y.imag(); + return __x; + } + template + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & + operator-=(complex &__x, const complex<_Xp> &__y) { + __x.__re_ -= __y.real(); + __x.__im_ -= __y.imag(); + return __x; + } + template + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & + operator*=(complex &__x, const complex<_Xp> &__y) { + __x = __x * complex(__y.real(), __y.imag()); + return __x; + } + template + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & + operator/=(complex &__x, const complex<_Xp> &__y) { + __x = __x / complex(__y.real(), __y.imag()); + return __x; + } + + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator+(const complex &__x, const complex &__y) { + complex __t(__x); + __t += __y; + return __t; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator+(const complex &__x, value_type __y) { + complex __t(__x); + __t += __y; + return __t; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator+(value_type __x, const complex &__y) { + complex __t(__y); + __t += __x; + return __t; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator+(const complex &__x) { + return __x; + } + + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator-(const complex &__x, const complex &__y) { + complex __t(__x); + __t -= __y; + return __t; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator-(const complex &__x, value_type __y) { + complex __t(__x); + __t -= __y; + return __t; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator-(value_type __x, const complex &__y) { + complex __t(-__y); + __t += __x; + return __t; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator-(const complex &__x) { + return complex(-__x.__re_, -__x.__im_); + } + + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator*(const complex &__z, const complex &__w) { + value_type __a = __z.__re_; + value_type __b = __z.__im_; + value_type __c = __w.__re_; + value_type __d = __w.__im_; + value_type __ac = __a * __c; + value_type __bd = __b * __d; + value_type __ad = __a * __d; + value_type __bc = __b * __c; + value_type __x = __ac - __bd; + value_type __y = __ad + __bc; + if (sycl::isnan(__x) && sycl::isnan(__y)) { + bool __recalc = false; + if (sycl::isinf(__a) || sycl::isinf(__b)) { + __a = sycl::copysign(sycl::isinf(__a) ? value_type(1) : value_type(0), + __a); + __b = sycl::copysign(sycl::isinf(__b) ? value_type(1) : value_type(0), + __b); + if (sycl::isnan(__c)) + __c = sycl::copysign(value_type(0), __c); + if (sycl::isnan(__d)) + __d = sycl::copysign(value_type(0), __d); + __recalc = true; + } + if (sycl::isinf(__c) || sycl::isinf(__d)) { + __c = sycl::copysign(sycl::isinf(__c) ? value_type(1) : value_type(0), + __c); + __d = sycl::copysign(sycl::isinf(__d) ? value_type(1) : value_type(0), + __d); + if (sycl::isnan(__a)) + __a = sycl::copysign(value_type(0), __a); + if (sycl::isnan(__b)) + __b = sycl::copysign(value_type(0), __b); + __recalc = true; + } + if (!__recalc && (sycl::isinf(__ac) || sycl::isinf(__bd) || + sycl::isinf(__ad) || sycl::isinf(__bc))) { + if (sycl::isnan(__a)) + __a = sycl::copysign(value_type(0), __a); + if (sycl::isnan(__b)) + __b = sycl::copysign(value_type(0), __b); + if (sycl::isnan(__c)) + __c = sycl::copysign(value_type(0), __c); + if (sycl::isnan(__d)) + __d = sycl::copysign(value_type(0), __d); + __recalc = true; + } + if (__recalc) { + __x = value_type(INFINITY) * (__a * __c - __b * __d); + __y = value_type(INFINITY) * (__a * __d + __b * __c); + } + } + return complex(__x, __y); + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator*(const complex &__x, value_type __y) { + complex __t(__x); + __t *= __y; + return __t; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator*(value_type __x, const complex &__y) { + complex __t(__y); + __t *= __x; + return __t; + } + + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator/(const complex &__z, const complex &__w) { + int __ilogbw = 0; + value_type __a = __z.__re_; + value_type __b = __z.__im_; + value_type __c = __w.__re_; + value_type __d = __w.__im_; + value_type __logbw = + sycl::logb(sycl::fmax(sycl::fabs(__c), sycl::fabs(__d))); + if (sycl::isfinite(__logbw)) { + __ilogbw = static_cast(__logbw); + __c = sycl::ldexp(__c, -__ilogbw); + __d = sycl::ldexp(__d, -__ilogbw); + } + value_type __denom = __c * __c + __d * __d; + value_type __x = sycl::ldexp((__a * __c + __b * __d) / __denom, -__ilogbw); + value_type __y = sycl::ldexp((__b * __c - __a * __d) / __denom, -__ilogbw); + if (sycl::isnan(__x) && sycl::isnan(__y)) { + if ((__denom == value_type(0)) && + (!sycl::isnan(__a) || !sycl::isnan(__b))) { + __x = sycl::copysign(value_type(INFINITY), __c) * __a; + __y = sycl::copysign(value_type(INFINITY), __c) * __b; + } else if ((sycl::isinf(__a) || sycl::isinf(__b)) && + sycl::isfinite(__c) && sycl::isfinite(__d)) { + __a = sycl::copysign(sycl::isinf(__a) ? value_type(1) : value_type(0), + __a); + __b = sycl::copysign(sycl::isinf(__b) ? value_type(1) : value_type(0), + __b); + __x = value_type(INFINITY) * (__a * __c + __b * __d); + __y = value_type(INFINITY) * (__b * __c - __a * __d); + } else if (sycl::isinf(__logbw) && __logbw > value_type(0) && + sycl::isfinite(__a) && sycl::isfinite(__b)) { + __c = sycl::copysign(sycl::isinf(__c) ? value_type(1) : value_type(0), + __c); + __d = sycl::copysign(sycl::isinf(__d) ? value_type(1) : value_type(0), + __d); + __x = value_type(0) * (__a * __c + __b * __d); + __y = value_type(0) * (__b * __c - __a * __d); + } + } + return complex(__x, __y); + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator/(const complex &__x, value_type __y) { + return complex(__x.__re_ / __y, __x.__im_ / __y); + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex + operator/(value_type __x, const complex &__y) { + complex __t(__x); + __t /= __y; + return __t; + } + + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend constexpr bool + operator==(const complex &__x, const complex &__y) { + return __x.__re_ == __y.__re_ && __x.__im_ == __y.__im_; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend constexpr bool + operator==(const complex &__x, value_type __y) { + return __x.__re_ == __y && __x.__im_ == 0; + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend constexpr bool + operator==(value_type __x, const complex &__y) { + return __x == __y.__re_ && 0 == __y.__im_; + } + + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend constexpr bool + operator!=(const complex &__x, const complex &__y) { + return !(__x == __y); + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend constexpr bool + operator!=(const complex &__x, value_type __y) { + return !(__x == __y); + } + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend constexpr bool + operator!=(value_type __x, const complex &__y) { + return !(__x == __y); + } + + template + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend std::basic_istream<_CharT, _Traits> & + operator>>(std::basic_istream<_CharT, _Traits> &__is, + complex &__x) { + if (__is.good()) { + ws(__is); + if (__is.peek() == _CharT('(')) { + __is.get(); + value_type __r; + __is >> __r; + if (!__is.fail()) { + ws(__is); + _CharT __c = __is.peek(); + if (__c == _CharT(',')) { + __is.get(); + value_type __i; + __is >> __i; + if (!__is.fail()) { + ws(__is); + __c = __is.peek(); + if (__c == _CharT(')')) { + __is.get(); + __x = complex(__r, __i); + } else + __is.setstate(__is.failbit); + } else + __is.setstate(__is.failbit); + } else if (__c == _CharT(')')) { + __is.get(); + __x = complex(__r, value_type(0)); + } else + __is.setstate(__is.failbit); + } else + __is.setstate(__is.failbit); + } else { + value_type __r; + __is >> __r; + if (!__is.fail()) + __x = complex(__r, value_type(0)); + else + __is.setstate(__is.failbit); + } + } else + __is.setstate(__is.failbit); + return __is; + } + + template + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend std::basic_ostream<_CharT, _Traits> & + operator<<(std::basic_ostream<_CharT, _Traits> &__os, + const complex &__x) { + std::basic_ostringstream<_CharT, _Traits> __s; + __s.flags(__os.flags()); + __s.imbue(__os.getloc()); + __s.precision(__os.precision()); + __s << '(' << __x.__re_ << ',' << __x.__im_ << ')'; + return __os << __s.str(); + } + + _SYCL_EXT_CPLX_INLINE_VISIBILITY friend const sycl::stream & + operator<<(const sycl::stream &__ss, const complex &_x) { + return __ss << "(" << _x.__re_ << "," << _x.__im_ << ")"; + } +}; + +} // namespace experimental +} // namespace oneapi +} // namespace ext + +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/ext/oneapi/experimental/sycl_complex.hpp b/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex_math.hpp similarity index 58% rename from sycl/include/sycl/ext/oneapi/experimental/sycl_complex.hpp rename to sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex_math.hpp index 6b6ddfb36e1da..cca76b7e988a5 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/sycl_complex.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/complex/detail/complex_math.hpp @@ -1,4 +1,4 @@ -//===- sycl_complex.hpp ---------------------------------------------------===// +//===- complex_math.hpp ---------------------------------------------------===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -8,37 +8,26 @@ #pragma once -#ifdef SYCL_EXT_ONEAPI_COMPLEX +#include "common.hpp" -#define _SYCL_EXT_CPLX_INLINE_VISIBILITY \ - inline __attribute__((__visibility__("hidden"), __always_inline__)) +#include -#include -#include // for std::basic_ostringstream -#include -#include +#include namespace sycl { inline namespace _V1 { + namespace ext { namespace oneapi { namespace experimental { -using std::enable_if; -using std::integral_constant; -using std::is_floating_point; -using std::is_integral; -using std::is_same; - -using std::basic_istream; -using std::basic_ostream; -using std::basic_ostringstream; - -using std::declval; +//////////////////////////////////////////////////////////////////////////////// +/// TRAITS +//////////////////////////////////////////////////////////////////////////////// namespace cplx::detail { -template using _BoolConstant = integral_constant; +template using _BoolConstant = std::integral_constant; template using _IsNotSame = _BoolConstant; @@ -56,7 +45,7 @@ template struct __numeric_type { static double __test(unsigned long long); static double __test(double); - typedef decltype(__test(declval<_Tp>())) type; + typedef decltype(__test(std::declval<_Tp>())) type; static const bool value = _IsNotSame::value; }; @@ -65,8 +54,8 @@ template <> struct __numeric_type { }; template ::value &&__numeric_type<_A2>::value - &&__numeric_type<_A3>::value> + bool = __numeric_type<_A1>::value && __numeric_type<_A2>::value && + __numeric_type<_A3>::value> class __promote_imp { public: static const bool value = false; @@ -102,383 +91,7 @@ template class __promote_imp<_A1, void, void, true> { template class __promote : public __promote_imp<_A1, _A2, _A3> {}; -} // namespace cplx::detail - -template class complex; - -template -struct is_gencomplex - : std::integral_constant> || - std::is_same_v<_Tp, complex> || - std::is_same_v<_Tp, complex>> {}; - -template -struct is_genfloat - : std::integral_constant || - std::is_same_v<_Tp, float> || - std::is_same_v<_Tp, sycl::half>> {}; - -template -class complex<_Tp, typename std::enable_if_t::value>> { -public: - typedef _Tp value_type; - -private: - value_type __re_; - value_type __im_; - -public: - _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr complex( - value_type __re = value_type(), value_type __im = value_type()) - : __re_(__re), __im_(__im) {} - - template - _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr complex(const complex<_Xp> &__c) - : __re_(__c.real()), __im_(__c.imag()) {} - - template ::value>> - _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr complex( - const std::complex<_Xp> &__c) - : __re_(static_cast(__c.real())), - __im_(static_cast(__c.imag())) {} - - template ::value>> - _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr - operator std::complex<_Xp>() const { - return std::complex<_Xp>(static_cast<_Xp>(__re_), static_cast<_Xp>(__im_)); - } - - _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr value_type real() const { - return __re_; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr value_type imag() const { - return __im_; - } - - _SYCL_EXT_CPLX_INLINE_VISIBILITY void real(value_type __re) { __re_ = __re; } - _SYCL_EXT_CPLX_INLINE_VISIBILITY void imag(value_type __im) { __im_ = __im; } - - _SYCL_EXT_CPLX_INLINE_VISIBILITY complex &operator=(value_type __re) { - __re_ = __re; - __im_ = value_type(); - return *this; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & - operator+=(complex &__c, value_type __re) { - __c.__re_ += __re; - return __c; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & - operator-=(complex &__c, value_type __re) { - __c.__re_ -= __re; - return __c; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & - operator*=(complex &__c, value_type __re) { - __c.__re_ *= __re; - __c.__im_ *= __re; - return __c; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & - operator/=(complex &__c, value_type __re) { - __c.__re_ /= __re; - __c.__im_ /= __re; - return __c; - } - - template - _SYCL_EXT_CPLX_INLINE_VISIBILITY complex &operator=(const complex<_Xp> &__c) { - __re_ = __c.real(); - __im_ = __c.imag(); - return *this; - } - template - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & - operator+=(complex &__x, const complex<_Xp> &__y) { - __x.__re_ += __y.real(); - __x.__im_ += __y.imag(); - return __x; - } - template - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & - operator-=(complex &__x, const complex<_Xp> &__y) { - __x.__re_ -= __y.real(); - __x.__im_ -= __y.imag(); - return __x; - } - template - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & - operator*=(complex &__x, const complex<_Xp> &__y) { - __x = __x * complex(__y.real(), __y.imag()); - return __x; - } - template - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex & - operator/=(complex &__x, const complex<_Xp> &__y) { - __x = __x / complex(__y.real(), __y.imag()); - return __x; - } - - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator+(const complex &__x, const complex &__y) { - complex __t(__x); - __t += __y; - return __t; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator+(const complex &__x, value_type __y) { - complex __t(__x); - __t += __y; - return __t; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator+(value_type __x, const complex &__y) { - complex __t(__y); - __t += __x; - return __t; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator+(const complex &__x) { - return __x; - } - - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator-(const complex &__x, const complex &__y) { - complex __t(__x); - __t -= __y; - return __t; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator-(const complex &__x, value_type __y) { - complex __t(__x); - __t -= __y; - return __t; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator-(value_type __x, const complex &__y) { - complex __t(-__y); - __t += __x; - return __t; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator-(const complex &__x) { - return complex(-__x.__re_, -__x.__im_); - } - - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator*(const complex &__z, const complex &__w) { - value_type __a = __z.__re_; - value_type __b = __z.__im_; - value_type __c = __w.__re_; - value_type __d = __w.__im_; - value_type __ac = __a * __c; - value_type __bd = __b * __d; - value_type __ad = __a * __d; - value_type __bc = __b * __c; - value_type __x = __ac - __bd; - value_type __y = __ad + __bc; - if (sycl::isnan(__x) && sycl::isnan(__y)) { - bool __recalc = false; - if (sycl::isinf(__a) || sycl::isinf(__b)) { - __a = sycl::copysign(sycl::isinf(__a) ? value_type(1) : value_type(0), - __a); - __b = sycl::copysign(sycl::isinf(__b) ? value_type(1) : value_type(0), - __b); - if (sycl::isnan(__c)) - __c = sycl::copysign(value_type(0), __c); - if (sycl::isnan(__d)) - __d = sycl::copysign(value_type(0), __d); - __recalc = true; - } - if (sycl::isinf(__c) || sycl::isinf(__d)) { - __c = sycl::copysign(sycl::isinf(__c) ? value_type(1) : value_type(0), - __c); - __d = sycl::copysign(sycl::isinf(__d) ? value_type(1) : value_type(0), - __d); - if (sycl::isnan(__a)) - __a = sycl::copysign(value_type(0), __a); - if (sycl::isnan(__b)) - __b = sycl::copysign(value_type(0), __b); - __recalc = true; - } - if (!__recalc && (sycl::isinf(__ac) || sycl::isinf(__bd) || - sycl::isinf(__ad) || sycl::isinf(__bc))) { - if (sycl::isnan(__a)) - __a = sycl::copysign(value_type(0), __a); - if (sycl::isnan(__b)) - __b = sycl::copysign(value_type(0), __b); - if (sycl::isnan(__c)) - __c = sycl::copysign(value_type(0), __c); - if (sycl::isnan(__d)) - __d = sycl::copysign(value_type(0), __d); - __recalc = true; - } - if (__recalc) { - __x = value_type(INFINITY) * (__a * __c - __b * __d); - __y = value_type(INFINITY) * (__a * __d + __b * __c); - } - } - return complex(__x, __y); - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator*(const complex &__x, value_type __y) { - complex __t(__x); - __t *= __y; - return __t; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator*(value_type __x, const complex &__y) { - complex __t(__y); - __t *= __x; - return __t; - } - - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator/(const complex &__z, const complex &__w) { - int __ilogbw = 0; - value_type __a = __z.__re_; - value_type __b = __z.__im_; - value_type __c = __w.__re_; - value_type __d = __w.__im_; - value_type __logbw = - sycl::logb(sycl::fmax(sycl::fabs(__c), sycl::fabs(__d))); - if (sycl::isfinite(__logbw)) { - __ilogbw = static_cast(__logbw); - __c = sycl::ldexp(__c, -__ilogbw); - __d = sycl::ldexp(__d, -__ilogbw); - } - value_type __denom = __c * __c + __d * __d; - value_type __x = sycl::ldexp((__a * __c + __b * __d) / __denom, -__ilogbw); - value_type __y = sycl::ldexp((__b * __c - __a * __d) / __denom, -__ilogbw); - if (sycl::isnan(__x) && sycl::isnan(__y)) { - if ((__denom == value_type(0)) && - (!sycl::isnan(__a) || !sycl::isnan(__b))) { - __x = sycl::copysign(value_type(INFINITY), __c) * __a; - __y = sycl::copysign(value_type(INFINITY), __c) * __b; - } else if ((sycl::isinf(__a) || sycl::isinf(__b)) && - sycl::isfinite(__c) && sycl::isfinite(__d)) { - __a = sycl::copysign(sycl::isinf(__a) ? value_type(1) : value_type(0), - __a); - __b = sycl::copysign(sycl::isinf(__b) ? value_type(1) : value_type(0), - __b); - __x = value_type(INFINITY) * (__a * __c + __b * __d); - __y = value_type(INFINITY) * (__b * __c - __a * __d); - } else if (sycl::isinf(__logbw) && __logbw > value_type(0) && - sycl::isfinite(__a) && sycl::isfinite(__b)) { - __c = sycl::copysign(sycl::isinf(__c) ? value_type(1) : value_type(0), - __c); - __d = sycl::copysign(sycl::isinf(__d) ? value_type(1) : value_type(0), - __d); - __x = value_type(0) * (__a * __c + __b * __d); - __y = value_type(0) * (__b * __c - __a * __d); - } - } - return complex(__x, __y); - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator/(const complex &__x, value_type __y) { - return complex(__x.__re_ / __y, __x.__im_ / __y); - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend complex - operator/(value_type __x, const complex &__y) { - complex __t(__x); - __t /= __y; - return __t; - } - - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend constexpr bool - operator==(const complex &__x, const complex &__y) { - return __x.__re_ == __y.__re_ && __x.__im_ == __y.__im_; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend constexpr bool - operator==(const complex &__x, value_type __y) { - return __x.__re_ == __y && __x.__im_ == 0; - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend constexpr bool - operator==(value_type __x, const complex &__y) { - return __x == __y.__re_ && 0 == __y.__im_; - } - - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend constexpr bool - operator!=(const complex &__x, const complex &__y) { - return !(__x == __y); - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend constexpr bool - operator!=(const complex &__x, value_type __y) { - return !(__x == __y); - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend constexpr bool - operator!=(value_type __x, const complex &__y) { - return !(__x == __y); - } - - template - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend std::basic_istream<_CharT, _Traits> & - operator>>(std::basic_istream<_CharT, _Traits> &__is, - complex &__x) { - if (__is.good()) { - ws(__is); - if (__is.peek() == _CharT('(')) { - __is.get(); - value_type __r; - __is >> __r; - if (!__is.fail()) { - ws(__is); - _CharT __c = __is.peek(); - if (__c == _CharT(',')) { - __is.get(); - value_type __i; - __is >> __i; - if (!__is.fail()) { - ws(__is); - __c = __is.peek(); - if (__c == _CharT(')')) { - __is.get(); - __x = complex(__r, __i); - } else - __is.setstate(__is.failbit); - } else - __is.setstate(__is.failbit); - } else if (__c == _CharT(')')) { - __is.get(); - __x = complex(__r, value_type(0)); - } else - __is.setstate(__is.failbit); - } else - __is.setstate(__is.failbit); - } else { - value_type __r; - __is >> __r; - if (!__is.fail()) - __x = complex(__r, value_type(0)); - else - __is.setstate(__is.failbit); - } - } else - __is.setstate(__is.failbit); - return __is; - } - - template - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend std::basic_ostream<_CharT, _Traits> & - operator<<(std::basic_ostream<_CharT, _Traits> &__os, - const complex &__x) { - std::basic_ostringstream<_CharT, _Traits> __s; - __s.flags(__os.flags()); - __s.imbue(__os.getloc()); - __s.precision(__os.precision()); - __s << '(' << __x.__re_ << ',' << __x.__im_ << ')'; - return __os << __s.str(); - } - _SYCL_EXT_CPLX_INLINE_VISIBILITY friend const sycl::stream & - operator<<(const sycl::stream &__ss, const complex &_x) { - return __ss << "(" << _x.__re_ << "," << _x.__im_ << ")"; - } -}; - -namespace cplx::detail { template ::value, bool = is_genfloat<_Tp>::value> struct __libcpp_complex_overload_traits {}; @@ -494,39 +107,12 @@ template struct __libcpp_complex_overload_traits<_Tp, false, true> { typedef _Tp _ValueType; typedef complex<_Tp> _ComplexType; }; -} // namespace cplx::detail - -// real - -template -__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr - typename std::enable_if_t::value, _Tp> - real(const complex<_Tp> &__c) { - return __c.real(); -} -template -__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr - typename cplx::detail::__libcpp_complex_overload_traits<_Tp>::_ValueType - real(_Tp __re) { - return __re; -} - -// imag - -template -__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr - typename std::enable_if_t::value, _Tp> - imag(const complex<_Tp> &__c) { - return __c.imag(); -} +} // namespace cplx::detail -template -__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr - typename cplx::detail::__libcpp_complex_overload_traits<_Tp>::_ValueType - imag(_Tp) { - return 0; -} +//////////////////////////////////////////////////////////////////////////////// +/// FUNCTIONS +//////////////////////////////////////////////////////////////////////////////// // abs @@ -1006,13 +592,41 @@ __DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY return complex<_Tp>(__z.imag(), -__z.real()); } +// real + +template +__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr + typename std::enable_if_t::value, _Tp> + real(const complex<_Tp> &__c) { + return __c.real(); +} + +template +__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr + typename cplx::detail::__libcpp_complex_overload_traits<_Tp>::_ValueType + real(_Tp __re) { + return __re; +} + +// imag + +template +__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr + typename std::enable_if_t::value, _Tp> + imag(const complex<_Tp> &__c) { + return __c.imag(); +} + +template +__DPCPP_SYCL_EXTERNAL _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr + typename cplx::detail::__libcpp_complex_overload_traits<_Tp>::_ValueType + imag(_Tp) { + return 0; +} + } // namespace experimental } // namespace oneapi } // namespace ext } // namespace _V1 } // namespace sycl - -#undef _SYCL_EXT_CPLX_INLINE_VISIBILITY - -#endif // SYCL_EXT_ONEAPI_COMPLEX diff --git a/sycl/include/syclcompat/util.hpp b/sycl/include/syclcompat/util.hpp index 848eaaefcd371..5978318340c08 100644 --- a/sycl/include/syclcompat/util.hpp +++ b/sycl/include/syclcompat/util.hpp @@ -35,7 +35,7 @@ #include #include -#include +#include #include #include diff --git a/sycl/test-e2e/Complex/sycl_complex_helper.hpp b/sycl/test-e2e/Complex/sycl_complex_helper.hpp index de4cc1813abf4..a57f736f5244e 100644 --- a/sycl/test-e2e/Complex/sycl_complex_helper.hpp +++ b/sycl/test-e2e/Complex/sycl_complex_helper.hpp @@ -2,7 +2,7 @@ #include #define SYCL_EXT_ONEAPI_COMPLEX -#include +#include #include using namespace sycl::ext::oneapi; diff --git a/sycl/test/basic_tests/macros_no_rdc.cpp b/sycl/test/basic_tests/macros_no_rdc.cpp index e5d0726847fa1..4ffc0ec8d91ca 100644 --- a/sycl/test/basic_tests/macros_no_rdc.cpp +++ b/sycl/test/basic_tests/macros_no_rdc.cpp @@ -25,4 +25,4 @@ #include #include "ext/oneapi/bfloat16.hpp" #include "ext/intel/esimd.hpp" -#include "ext/oneapi/experimental/sycl_complex.hpp" +#include "ext/oneapi/experimental/complex/complex.hpp" diff --git a/sycl/test/extensions/test_complex.cpp b/sycl/test/extensions/complex/complex.cpp similarity index 95% rename from sycl/test/extensions/test_complex.cpp rename to sycl/test/extensions/complex/complex.cpp index 60b466901e0d8..de39f9f8472cd 100644 --- a/sycl/test/extensions/test_complex.cpp +++ b/sycl/test/extensions/complex/complex.cpp @@ -2,7 +2,7 @@ #define SYCL_EXT_ONEAPI_COMPLEX -#include +#include #include using namespace sycl::ext::oneapi::experimental; @@ -24,8 +24,9 @@ template