Skip to content

Commit

Permalink
Merge pull request IntelPython#1707 from IntelPython/replace-some-std…
Browse files Browse the repository at this point in the history
…-namepsace-functions-with-sycl-namespace-functions

Remove USE_SYCL_FOR_COMPLEX_TYPES
  • Loading branch information
ndgrigorian authored Jun 10, 2024
2 parents f114571 + 6ea1120 commit c81735a
Show file tree
Hide file tree
Showing 38 changed files with 124 additions and 491 deletions.
2 changes: 1 addition & 1 deletion dpctl/tensor/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -231,7 +231,7 @@ foreach(_src_fn ${_no_fast_math_sources})
)
endforeach()

set(_compiler_definitions "USE_SYCL_FOR_COMPLEX_TYPES")
set(_compiler_definitions "")

foreach(_src_fn ${_elementwise_sources})
get_source_file_property(_cmpl_options_defs ${_src_fn} COMPILE_DEFINITIONS)
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -78,7 +78,7 @@ template <typename argT, typename resT> struct AbsFunctor
else if constexpr (std::is_same_v<argT, sycl::half> ||
std::is_floating_point_v<argT>)
{
return (std::signbit(x) ? -x : x);
return (sycl::signbit(x) ? -x : x);
}
else {
return sycl::abs(x);
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -89,7 +89,7 @@ template <typename argT, typename resT> struct AcosFunctor
}
/* acos(0 + I*NaN) = PI/2 + I*NaN with inexact */
if (x == realT(0)) {
const realT res_re = std::atan(realT(1)) * 2; // PI/2
const realT res_re = sycl::atan(realT(1)) * 2; // PI/2
return resT{res_re, q_nan};
}

Expand All @@ -103,7 +103,6 @@ template <typename argT, typename resT> struct AcosFunctor
constexpr realT r_eps =
realT(1) / std::numeric_limits<realT>::epsilon();
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
using sycl_complexT = exprm_ns::complex<realT>;
sycl_complexT log_in =
exprm_ns::log(exprm_ns::complex<realT>(in));
Expand All @@ -112,31 +111,17 @@ template <typename argT, typename resT> struct AcosFunctor
const realT wy = log_in.imag();
const realT rx = sycl::fabs(wy);

realT ry = wx + std::log(realT(2));
return resT{rx, (std::signbit(y)) ? ry : -ry};
#else
resT log_in = std::log(in);
const realT wx = std::real(log_in);
const realT wy = std::imag(log_in);
const realT rx = sycl::fabs(wy);

realT ry = wx + std::log(realT(2));
return resT{rx, (std::signbit(y)) ? ry : -ry};
#endif
realT ry = wx + sycl::log(realT(2));
return resT{rx, (sycl::signbit(y)) ? ry : -ry};
}

/* ordinary cases */
#if USE_SYCL_FOR_COMPLEX_TYPES
return exprm_ns::acos(
exprm_ns::complex<realT>(in)); // std::acos(in);
#else
return std::acos(in);
#endif
return exprm_ns::acos(exprm_ns::complex<realT>(in)); // acos(in);
}
else {
static_assert(std::is_floating_point_v<argT> ||
std::is_same_v<argT, sycl::half>);
return std::acos(in);
return sycl::acos(in);
}
}
};
Expand Down Expand Up @@ -203,7 +188,7 @@ template <typename fnT, typename T> struct AcosContigFactory

template <typename fnT, typename T> struct AcosTypeMapFactory
{
/*! @brief get typeid for output type of std::acos(T x) */
/*! @brief get typeid for output type of sycl::acos(T x) */
std::enable_if_t<std::is_same<fnT, int>::value, int> get()
{
using rT = typename AcosOutputType<T>::value_type;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -96,7 +96,7 @@ template <typename argT, typename resT> struct AcoshFunctor
}
/* acos(0 + I*NaN) = Pi/2 + I*NaN with inexact */
else if (x == realT(0)) {
const realT pi_half = std::atan(1) * 2;
const realT pi_half = sycl::atan(realT(1)) * 2;
acos_in = resT{pi_half, q_nan};
}
else {
Expand All @@ -110,28 +110,18 @@ template <typename argT, typename resT> struct AcoshFunctor
* For large x or y including acos(+-Inf + I*+-Inf)
*/
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
using sycl_complexT = typename exprm_ns::complex<realT>;
const sycl_complexT log_in = exprm_ns::log(sycl_complexT(in));
const realT wx = log_in.real();
const realT wy = log_in.imag();
#else
const resT log_in = std::log(in);
const realT wx = std::real(log_in);
const realT wy = std::imag(log_in);
#endif
const realT rx = sycl::fabs(wy);
realT ry = wx + std::log(realT(2));
acos_in = resT{rx, (std::signbit(y)) ? ry : -ry};
realT ry = wx + sycl::log(realT(2));
acos_in = resT{rx, (sycl::signbit(y)) ? ry : -ry};
}
else {
/* ordinary cases */
#if USE_SYCL_FOR_COMPLEX_TYPES
acos_in = exprm_ns::acos(
exprm_ns::complex<realT>(in)); // std::acos(in);
#else
acos_in = std::acos(in);
#endif
acos_in =
exprm_ns::acos(exprm_ns::complex<realT>(in)); // acos(in);
}

/* Now we calculate acosh(z) */
Expand All @@ -158,7 +148,7 @@ template <typename argT, typename resT> struct AcoshFunctor
else {
static_assert(std::is_floating_point_v<argT> ||
std::is_same_v<argT, sycl::half>);
return std::acosh(in);
return sycl::acosh(in);
}
}
};
Expand Down Expand Up @@ -225,7 +215,7 @@ template <typename fnT, typename T> struct AcoshContigFactory

template <typename fnT, typename T> struct AcoshTypeMapFactory
{
/*! @brief get typeid for output type of std::acosh(T x) */
/*! @brief get typeid for output type of sycl::acosh(T x) */
std::enable_if_t<std::is_same<fnT, int>::value, int> get()
{
using rT = typename AcoshOutputType<T>::value_type;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -63,36 +63,24 @@ template <typename argT1, typename argT2, typename resT> struct AddFunctor
if constexpr (tu_ns::is_complex<argT1>::value &&
tu_ns::is_complex<argT2>::value)
{
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
using rT1 = typename argT1::value_type;
using rT2 = typename argT2::value_type;

return exprm_ns::complex<rT1>(in1) + exprm_ns::complex<rT2>(in2);
#else
return in1 + in2;
#endif
}
else if constexpr (tu_ns::is_complex<argT1>::value &&
!tu_ns::is_complex<argT2>::value)
{
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
using rT1 = typename argT1::value_type;

return exprm_ns::complex<rT1>(in1) + in2;
#else
return in1 + in2;
#endif
}
else if constexpr (!tu_ns::is_complex<argT1>::value &&
tu_ns::is_complex<argT2>::value)
{
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
using rT2 = typename argT2::value_type;

return in1 + exprm_ns::complex<rT2>(in2);
#else
return in1 + in2;
#endif
}
else {
return in1 + in2;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -66,13 +66,9 @@ template <typename argT, typename resT> struct AngleFunctor

resT operator()(const argT &in) const
{
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
using rT = typename argT::value_type;

return exprm_ns::arg(exprm_ns::complex<rT>(in)); // std::arg(in);
#else
return std::arg(in);
#endif
return exprm_ns::arg(exprm_ns::complex<rT>(in)); // arg(in);
}
};

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -117,50 +117,31 @@ template <typename argT, typename resT> struct AsinFunctor
constexpr realT r_eps =
realT(1) / std::numeric_limits<realT>::epsilon();
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
using sycl_complexT = exprm_ns::complex<realT>;
const sycl_complexT z{x, y};
realT wx, wy;
if (!std::signbit(x)) {
if (!sycl::signbit(x)) {
const auto log_z = exprm_ns::log(z);
wx = log_z.real() + std::log(realT(2));
wx = log_z.real() + sycl::log(realT(2));
wy = log_z.imag();
}
else {
const auto log_mz = exprm_ns::log(-z);
wx = log_mz.real() + std::log(realT(2));
wx = log_mz.real() + sycl::log(realT(2));
wy = log_mz.imag();
}
#else
const resT z{x, y};
realT wx, wy;
if (!std::signbit(x)) {
const auto log_z = std::log(z);
wx = std::real(log_z) + std::log(realT(2));
wy = std::imag(log_z);
}
else {
const auto log_mz = std::log(-z);
wx = std::real(log_mz) + std::log(realT(2));
wy = std::imag(log_mz);
}
#endif
const realT asinh_re = sycl::copysign(wx, x);
const realT asinh_im = sycl::copysign(wy, y);
return resT{asinh_im, asinh_re};
}
/* ordinary cases */
#if USE_SYCL_FOR_COMPLEX_TYPES
return exprm_ns::asin(
exprm_ns::complex<realT>(in)); // std::asin(in);
#else
return std::asin(in);
#endif
exprm_ns::complex<realT>(in)); // sycl::asin(in);
}
else {
static_assert(std::is_floating_point_v<argT> ||
std::is_same_v<argT, sycl::half>);
return std::asin(in);
return sycl::asin(in);
}
}
};
Expand Down Expand Up @@ -227,7 +208,7 @@ template <typename fnT, typename T> struct AsinContigFactory

template <typename fnT, typename T> struct AsinTypeMapFactory
{
/*! @brief get typeid for output type of std::asin(T x) */
/*! @brief get typeid for output type of sycl::asin(T x) */
std::enable_if_t<std::is_same<fnT, int>::value, int> get()
{
using rT = typename AsinOutputType<T>::value_type;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -106,35 +106,25 @@ template <typename argT, typename resT> struct AsinhFunctor
realT(1) / std::numeric_limits<realT>::epsilon();

if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
using sycl_complexT = exprm_ns::complex<realT>;
sycl_complexT log_in = (std::signbit(x))
sycl_complexT log_in = (sycl::signbit(x))
? exprm_ns::log(sycl_complexT(-in))
: exprm_ns::log(sycl_complexT(in));
realT wx = log_in.real() + std::log(realT(2));
realT wx = log_in.real() + sycl::log(realT(2));
realT wy = log_in.imag();
#else
auto log_in = std::log(std::signbit(x) ? -in : in);
realT wx = std::real(log_in) + std::log(realT(2));
realT wy = std::imag(log_in);
#endif

const realT res_re = sycl::copysign(wx, x);
const realT res_im = sycl::copysign(wy, y);
return resT{res_re, res_im};
}

/* ordinary cases */
#if USE_SYCL_FOR_COMPLEX_TYPES
return exprm_ns::asinh(
exprm_ns::complex<realT>(in)); // std::asinh(in);
#else
return std::asinh(in);
#endif
return exprm_ns::asinh(exprm_ns::complex<realT>(in)); // asinh(in);
}
else {
static_assert(std::is_floating_point_v<argT> ||
std::is_same_v<argT, sycl::half>);
return std::asinh(in);
return sycl::asinh(in);
}
}
};
Expand Down Expand Up @@ -201,7 +191,7 @@ template <typename fnT, typename T> struct AsinhContigFactory

template <typename fnT, typename T> struct AsinhTypeMapFactory
{
/*! @brief get typeid for output type of std::asinh(T x) */
/*! @brief get typeid for output type of sycl::asinh(T x) */
std::enable_if_t<std::is_same<fnT, int>::value, int> get()
{
using rT = typename AsinhOutputType<T>::value_type;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -82,7 +82,7 @@ template <typename argT, typename resT> struct AtanFunctor
if (std::isnan(x)) {
/* atanh(NaN + I*+-Inf) = sign(NaN)*0 + I*+-Pi/2 */
if (std::isinf(y)) {
const realT pi_half = std::atan(realT(1)) * 2;
const realT pi_half = sycl::atan(realT(1)) * 2;

const realT atanh_re = sycl::copysign(realT(0), x);
const realT atanh_im = sycl::copysign(pi_half, y);
Expand Down Expand Up @@ -119,24 +119,19 @@ template <typename argT, typename resT> struct AtanFunctor
constexpr realT r_eps =
realT(1) / std::numeric_limits<realT>::epsilon();
if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) {
const realT pi_half = std::atan(realT(1)) * 2;
const realT pi_half = sycl::atan(realT(1)) * 2;

const realT atanh_re = realT(0);
const realT atanh_im = sycl::copysign(pi_half, y);
return resT{atanh_im, atanh_re};
}
/* ordinary cases */
#ifdef USE_SYCL_FOR_COMPLEX_TYPES
return exprm_ns::atan(
exprm_ns::complex<realT>(in)); // std::atan(in);
#else
return std::atan(in);
#endif
return exprm_ns::atan(exprm_ns::complex<realT>(in)); // atan(in);
}
else {
static_assert(std::is_floating_point_v<argT> ||
std::is_same_v<argT, sycl::half>);
return std::atan(in);
return sycl::atan(in);
}
}
};
Expand Down Expand Up @@ -203,7 +198,7 @@ template <typename fnT, typename T> struct AtanContigFactory

template <typename fnT, typename T> struct AtanTypeMapFactory
{
/*! @brief get typeid for output type of std::atan(T x) */
/*! @brief get typeid for output type of sycl::atan(T x) */
std::enable_if_t<std::is_same<fnT, int>::value, int> get()
{
using rT = typename AtanOutputType<T>::value_type;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -56,12 +56,12 @@ template <typename argT1, typename argT2, typename resT> struct Atan2Functor

resT operator()(const argT1 &in1, const argT2 &in2) const
{
if (std::isinf(in2) && !std::signbit(in2)) {
if (std::isinf(in2) && !sycl::signbit(in2)) {
if (std::isfinite(in1)) {
return sycl::copysign(resT(0), in1);
}
}
return std::atan2(in1, in2);
return sycl::atan2(in1, in2);
}
};

Expand Down Expand Up @@ -145,7 +145,7 @@ template <typename fnT, typename T1, typename T2> struct Atan2ContigFactory

template <typename fnT, typename T1, typename T2> struct Atan2TypeMapFactory
{
/*! @brief get typeid for output type of std::hypot(T1 x, T2 y) */
/*! @brief get typeid for output type of sycl::atan2(T1 x, T2 y) */
std::enable_if_t<std::is_same<fnT, int>::value, int> get()
{
using rT = typename Atan2OutputType<T1, T2>::value_type;
Expand Down
Loading

0 comments on commit c81735a

Please sign in to comment.