From 591bc916a3c6a6491f314eb74309a9b7d32e9112 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Sat, 8 Jun 2024 14:56:06 -0500 Subject: [PATCH 1/6] Remove USE_SYCL_NAMESPACE_FOR_COMPLEX_TYPES Always use SYCL namespace functions. Also replaced some std namespace functions with corresponding sycl namespace functions. --- dpctl/tensor/CMakeLists.txt | 2 +- .../kernels/elementwise_functions/abs.hpp | 2 +- .../kernels/elementwise_functions/acos.hpp | 24 +-- .../kernels/elementwise_functions/acosh.hpp | 20 +- .../kernels/elementwise_functions/add.hpp | 12 -- .../kernels/elementwise_functions/angle.hpp | 6 +- .../kernels/elementwise_functions/asin.hpp | 27 +-- .../kernels/elementwise_functions/asinh.hpp | 19 +- .../kernels/elementwise_functions/atan.hpp | 12 +- .../kernels/elementwise_functions/atanh.hpp | 8 +- .../elementwise_functions/cabs_impl.hpp | 4 - .../kernels/elementwise_functions/conj.hpp | 6 +- .../kernels/elementwise_functions/cos.hpp | 10 +- .../kernels/elementwise_functions/cosh.hpp | 10 +- .../kernels/elementwise_functions/equal.hpp | 4 - .../kernels/elementwise_functions/exp.hpp | 12 +- .../kernels/elementwise_functions/exp2.hpp | 8 +- .../kernels/elementwise_functions/log.hpp | 8 +- .../kernels/elementwise_functions/log10.hpp | 8 +- .../kernels/elementwise_functions/log2.hpp | 11 +- .../elementwise_functions/multiply.hpp | 4 - .../kernels/elementwise_functions/pow.hpp | 12 +- .../elementwise_functions/reciprocal.hpp | 4 - .../kernels/elementwise_functions/sin.hpp | 10 +- .../kernels/elementwise_functions/sinh.hpp | 10 +- .../kernels/elementwise_functions/sqrt.hpp | 191 +----------------- .../kernels/elementwise_functions/square.hpp | 4 - .../elementwise_functions/sycl_complex.hpp | 5 +- .../kernels/elementwise_functions/tan.hpp | 10 +- .../kernels/elementwise_functions/tanh.hpp | 10 +- .../elementwise_functions/true_divide.hpp | 20 -- .../kernels/integer_advanced_indexing.hpp | 4 +- 32 files changed, 67 insertions(+), 430 deletions(-) diff --git a/dpctl/tensor/CMakeLists.txt b/dpctl/tensor/CMakeLists.txt index 6639c09f9b..9cfcb5d46c 100644 --- a/dpctl/tensor/CMakeLists.txt +++ b/dpctl/tensor/CMakeLists.txt @@ -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) diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp index 0eafef8fe3..5ae983803a 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp @@ -78,7 +78,7 @@ template struct AbsFunctor else if constexpr (std::is_same_v || std::is_floating_point_v) { - return (std::signbit(x) ? -x : x); + return (sycl::signbit(x) ? -x : x); } else { return sycl::abs(x); diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp index 87a0d3b92a..0befcd8b74 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp @@ -89,7 +89,7 @@ template 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}; } @@ -103,7 +103,6 @@ template struct AcosFunctor constexpr realT r_eps = realT(1) / std::numeric_limits::epsilon(); if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using sycl_complexT = exprm_ns::complex; sycl_complexT log_in = exprm_ns::log(exprm_ns::complex(in)); @@ -112,31 +111,18 @@ template 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(in)); // std::acos(in); -#else - return std::acos(in); -#endif + exprm_ns::complex(in)); // acos(in); } else { static_assert(std::is_floating_point_v || std::is_same_v); - return std::acos(in); + return sycl::acos(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp index 1dfde6a523..1c18d08d35 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp @@ -96,7 +96,7 @@ template 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 { @@ -110,28 +110,18 @@ template 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; 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(in)); // std::acos(in); -#else - acos_in = std::acos(in); -#endif + exprm_ns::complex(in)); // acos(in); } /* Now we calculate acosh(z) */ @@ -158,7 +148,7 @@ template struct AcoshFunctor else { static_assert(std::is_floating_point_v || std::is_same_v); - return std::acosh(in); + return sycl::acosh(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp index 41078f0eae..bbd7045f04 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp @@ -63,36 +63,24 @@ template struct AddFunctor if constexpr (tu_ns::is_complex::value && tu_ns::is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using rT1 = typename argT1::value_type; using rT2 = typename argT2::value_type; return exprm_ns::complex(in1) + exprm_ns::complex(in2); -#else - return in1 + in2; -#endif } else if constexpr (tu_ns::is_complex::value && !tu_ns::is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using rT1 = typename argT1::value_type; return exprm_ns::complex(in1) + in2; -#else - return in1 + in2; -#endif } else if constexpr (!tu_ns::is_complex::value && tu_ns::is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using rT2 = typename argT2::value_type; return in1 + exprm_ns::complex(in2); -#else - return in1 + in2; -#endif } else { return in1 + in2; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/angle.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/angle.hpp index 0d1fb9cbf7..e34159d8c1 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/angle.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/angle.hpp @@ -66,13 +66,9 @@ template 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(in)); // std::arg(in); -#else - return std::arg(in); -#endif + return exprm_ns::arg(exprm_ns::complex(in)); // arg(in); } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp index f8504f8c8e..49f0513bf0 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp @@ -117,50 +117,31 @@ template struct AsinFunctor constexpr realT r_eps = realT(1) / std::numeric_limits::epsilon(); if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using sycl_complexT = exprm_ns::complex; 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(in)); // std::asin(in); -#else - return std::asin(in); -#endif } else { static_assert(std::is_floating_point_v || std::is_same_v); - return std::asin(in); + return sycl::asin(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp index df9a9a6e32..312dd6e43a 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp @@ -106,35 +106,26 @@ template struct AsinhFunctor realT(1) / std::numeric_limits::epsilon(); if (sycl::fabs(x) > r_eps || sycl::fabs(y) > r_eps) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using sycl_complexT = exprm_ns::complex; - 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(in)); // std::asinh(in); -#else - return std::asinh(in); -#endif + exprm_ns::complex(in)); // asinh(in); } else { static_assert(std::is_floating_point_v || std::is_same_v); - return std::asinh(in); + return sycl::asinh(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp index c03aed6e7f..f7b57c3487 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp @@ -82,7 +82,7 @@ template 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); @@ -119,24 +119,20 @@ template struct AtanFunctor constexpr realT r_eps = realT(1) / std::numeric_limits::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(in)); // std::atan(in); -#else - return std::atan(in); -#endif + exprm_ns::complex(in)); // atan(in); } else { static_assert(std::is_floating_point_v || std::is_same_v); - return std::atan(in); + return sycl::atan(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp index f4a33d0548..2fcee9ad14 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp @@ -120,17 +120,13 @@ template struct AtanhFunctor return resT{res_re, res_im}; } /* ordinary cases */ -#ifdef USE_SYCL_FOR_COMPLEX_TYPES return exprm_ns::atanh( - exprm_ns::complex(in)); // std::atanh(in); -#else - return std::atanh(in); -#endif + exprm_ns::complex(in)); // atanh(in); } else { static_assert(std::is_floating_point_v || std::is_same_v); - return std::atanh(in); + return sycl::atanh(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cabs_impl.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cabs_impl.hpp index 77930fcd35..e61304bed8 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cabs_impl.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cabs_impl.hpp @@ -70,11 +70,7 @@ template realT cabs(std::complex const &z) return q_nan; } else { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES return exprm_ns::abs(exprm_ns::complex(z)); -#else - return std::hypot(std::real(z), std::imag(z)); -#endif } } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp index 02963a5412..e0123e8d91 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp @@ -68,13 +68,9 @@ template struct ConjFunctor resT operator()(const argT &in) const { if constexpr (is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using rT = typename argT::value_type; - return exprm_ns::conj(exprm_ns::complex(in)); // std::conj(in); -#else - return std::conj(in); -#endif + return exprm_ns::conj(exprm_ns::complex(in)); // conj(in); } else { if constexpr (!std::is_same_v) diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp index 8aa8d8f18f..74ba53e40b 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp @@ -81,12 +81,8 @@ template struct CosFunctor * real and imaginary parts of input are finite. */ if (in_re_finite && in_im_finite) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES return exprm_ns::cos( - exprm_ns::complex(in)); // std::cos(in); -#else - return std::cos(in); -#endif + exprm_ns::complex(in)); // cos(in); } /* @@ -146,7 +142,7 @@ template struct CosFunctor if (!yfinite) { return resT{x * x, sycl::copysign(q_nan, x)}; } - return resT{(x * x) * std::cos(y), x * std::sin(y)}; + return resT{(x * x) * sycl::cos(y), x * sycl::sin(y)}; } /* @@ -161,7 +157,7 @@ template struct CosFunctor else { static_assert(std::is_floating_point_v || std::is_same_v); - return std::cos(in); + return sycl::cos(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp index c73d2caa61..2c3c2643cb 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp @@ -81,12 +81,8 @@ template struct CoshFunctor * real and imaginary parts of input are finite. */ if (xfinite && yfinite) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES return exprm_ns::cosh( - exprm_ns::complex(in)); // std::cosh(in); -#else - return std::cosh(in); -#endif + exprm_ns::complex(in)); // cosh(in); } /* @@ -135,7 +131,7 @@ template struct CoshFunctor if (!yfinite) { return resT{x * x, x * q_nan}; } - return resT{(x * x) * std::cos(y), x * std::sin(y)}; + return resT{(x * x) * sycl::cos(y), x * sycl::sin(y)}; } /* @@ -150,7 +146,7 @@ template struct CoshFunctor else { static_assert(std::is_floating_point_v || std::is_same_v); - return std::cosh(in); + return sycl::cosh(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/equal.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/equal.hpp index 61ac3ca128..b3a1627a9a 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/equal.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/equal.hpp @@ -68,12 +68,8 @@ template struct EqualFunctor using realT1 = typename argT1::value_type; using realT2 = typename argT2::value_type; -#ifdef USE_SYCL_FOR_COMPLEX_TYPES return exprm_ns::complex(in1) == exprm_ns::complex(in2); -#else - return (in1 == in2); -#endif } else { if constexpr (std::is_integral_v && diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp index aba24b81f6..fbca696a1a 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp @@ -73,12 +73,8 @@ template struct ExpFunctor const realT y = std::imag(in); if (std::isfinite(x)) { if (std::isfinite(y)) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES return exprm_ns::exp( - exprm_ns::complex(in)); // std::exp(in); -#else - return std::exp(in); -#endif + exprm_ns::complex(in)); // exp(in); } else { return resT{q_nan, q_nan}; @@ -99,7 +95,7 @@ template struct ExpFunctor return resT{x, y}; } else if (std::isfinite(y)) { - return resT{x * std::cos(y), x * std::sin(y)}; + return resT{x * sycl::cos(y), x * sycl::sin(y)}; } else { /* x = +inf, y = +-inf || nan */ @@ -109,7 +105,7 @@ template struct ExpFunctor else { /* x is -inf */ if (std::isfinite(y)) { realT exp_x = std::exp(x); - return resT{exp_x * std::cos(y), exp_x * std::sin(y)}; + return resT{exp_x * sycl::cos(y), exp_x * sycl::sin(y)}; } else { /* x = -inf, y = +-inf || nan */ @@ -119,7 +115,7 @@ template struct ExpFunctor } } else { - return std::exp(in); + return sycl::exp(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp index 2265bcce40..5bc652bd0d 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp @@ -76,11 +76,7 @@ template struct Exp2Functor const realT y = std::imag(tmp); if (std::isfinite(x)) { if (std::isfinite(y)) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES return exprm_ns::exp(exprm_ns::complex(tmp)); -#else - return std::exp(tmp); -#endif } else { return resT{q_nan, q_nan}; @@ -101,7 +97,7 @@ template struct Exp2Functor return resT{x, y}; } else if (std::isfinite(y)) { - return resT{x * std::cos(y), x * std::sin(y)}; + return resT{x * sycl::cos(y), x * sycl::sin(y)}; } else { /* x = +inf, y = +-inf || nan */ @@ -111,7 +107,7 @@ template struct Exp2Functor else { /* x is -inf */ if (std::isfinite(y)) { realT exp_x = std::exp(x); - return resT{exp_x * std::cos(y), exp_x * std::sin(y)}; + return resT{exp_x * sycl::cos(y), exp_x * sycl::sin(y)}; } else { /* x = -inf, y = +-inf || nan */ diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp index 321e97e8e0..fa2bc20b2f 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp @@ -67,15 +67,11 @@ template struct LogFunctor resT operator()(const argT &in) const { if constexpr (is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using realT = typename argT::value_type; - return exprm_ns::log(exprm_ns::complex(in)); // std::log(in); -#else - return std::log(in); -#endif + return exprm_ns::log(exprm_ns::complex(in)); // log(in); } else { - return std::log(in); + return sycl::log(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp index 7a18444a08..a7ce78a757 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp @@ -71,16 +71,12 @@ template struct Log10Functor { if constexpr (is_complex::value) { using realT = typename argT::value_type; -#ifdef USE_SYCL_FOR_COMPLEX_TYPES - // return (std::log(in) / std::log(realT{10})); + // return (log(in) / log(realT{10})); return exprm_ns::log(exprm_ns::complex(in)) / std::log(realT{10}); -#else - return (std::log(in) / std::log(realT{10})); -#endif } else { - return std::log10(in); + return sycl::log10(in); } } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp index ca39b2b8db..b98d93f6eb 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp @@ -71,16 +71,13 @@ template struct Log2Functor { if constexpr (is_complex::value) { using realT = typename argT::value_type; -#ifdef USE_SYCL_FOR_COMPLEX_TYPES - // std::log(in) / std::log(realT{2}); + + // log(in) / log(realT{2}); return exprm_ns::log(exprm_ns::complex(in)) / - std::log(realT{2}); -#else - return std::log(in) / std::log(realT{2}); -#endif + sycl::log(realT{2}); } else { - return std::log2(in); + return sycl::log2(in); } } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/multiply.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/multiply.hpp index 2c6448ec1a..b98ca2ef16 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/multiply.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/multiply.hpp @@ -64,15 +64,11 @@ template struct MultiplyFunctor if constexpr (tu_ns::is_complex::value && tu_ns::is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using realT1 = typename argT1::value_type; using realT2 = typename argT2::value_type; return exprm_ns::complex(in1) * exprm_ns::complex(in2); -#else - return in1 * in2; -#endif } else { return in1 * in2; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/pow.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/pow.hpp index 7ef9b302d4..7b2fe38af5 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/pow.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/pow.hpp @@ -86,18 +86,14 @@ template struct PowFunctor else if constexpr (tu_ns::is_complex::value && tu_ns::is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using realT1 = typename argT1::value_type; using realT2 = typename argT2::value_type; return exprm_ns::pow(exprm_ns::complex(in1), exprm_ns::complex(in2)); -#else - return std::pow(in1, in2); -#endif } else { - return std::pow(in1, in2); + return sycl::pow(in1, in2); } } @@ -369,18 +365,14 @@ template struct PowInplaceFunctor else if constexpr (tu_ns::is_complex::value && tu_ns::is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using r_resT = typename resT::value_type; using r_argT = typename argT::value_type; res = exprm_ns::pow(exprm_ns::complex(res), exprm_ns::complex(in)); -#else - res = std::pow(res, in); -#endif } else { - res = std::pow(res, in); + res = sycl::pow(res, in); } return; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/reciprocal.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/reciprocal.hpp index af9374a9da..76c48e173c 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/reciprocal.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/reciprocal.hpp @@ -70,12 +70,8 @@ template struct ReciprocalFunctor if constexpr (is_complex::value) { using realT = typename argT::value_type; -#ifdef USE_SYCL_FOR_COMPLEX_TYPES return realT(1) / exprm_ns::complex(in); -#else - return realT(1) / in; -#endif } else { return argT(1) / in; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp index e6570e4891..2bc38bfc12 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp @@ -79,12 +79,8 @@ template struct SinFunctor * real and imaginary parts of input are finite. */ if (in_re_finite && in_im_finite) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES resT res = exprm_ns::sin( - exprm_ns::complex(in)); // std::sin(in); -#else - resT res = std::sin(in); -#endif + exprm_ns::complex(in)); // sin(in); if (in_re == realT(0)) { res.real(sycl::copysign(realT(0), in_re)); } @@ -160,7 +156,7 @@ template struct SinFunctor const realT sinh_im = x * (y - y); return resT{sinh_im, -sinh_re}; } - const realT sinh_re = x * std::cos(y); + const realT sinh_re = x * sycl::cos(y); const realT sinh_im = std::numeric_limits::infinity() * std::sin(y); return resT{sinh_im, -sinh_re}; @@ -183,7 +179,7 @@ template struct SinFunctor if (in == 0) { return in; } - return std::sin(in); + return sycl::sin(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp index 436fc7ed82..64dbde7eef 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp @@ -79,11 +79,7 @@ template struct SinhFunctor * real and imaginary parts of input are finite. */ if (xfinite && yfinite) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES return exprm_ns::sinh(exprm_ns::complex(in)); -#else - return std::sinh(in); -#endif } /* * sinh(+-0 +- I Inf) = sign(d(+-0, dNaN))0 + I dNaN. @@ -135,9 +131,9 @@ template struct SinhFunctor if (!yfinite) { return resT{x * x, x * (y - y)}; } - return resT{x * std::cos(y), + return resT{x * sycl::cos(y), std::numeric_limits::infinity() * - std::sin(y)}; + sycl::sin(y)}; } /* @@ -152,7 +148,7 @@ template struct SinhFunctor else { static_assert(std::is_floating_point_v || std::is_same_v); - return std::sinh(in); + return sycl::sinh(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sqrt.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sqrt.hpp index 471f258fc2..032666618b 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sqrt.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sqrt.hpp @@ -70,201 +70,12 @@ template struct SqrtFunctor { if constexpr (is_complex::value) { using realT = typename argT::value_type; -#ifdef USE_SYCL_FOR_COMPLEX_TYPES return exprm_ns::sqrt(exprm_ns::complex(in)); -#else -#ifdef _WINDOWS - // Work around a problem on Windows, where std::sqrt for - // single precision input uses double type, precluding - // offloading to devices that do not support double precision - // i.e. Iris Xe - if constexpr (std::is_same_v) { - return std::sqrt(in); - } - else { - return csqrt(in); - } -#else - return std::sqrt(in); -#endif -#endif } else { - return std::sqrt(in); + return sycl::sqrt(in); } } - -private: - template std::complex csqrt(std::complex const &z) const - { - // csqrt(x + y*1j) - // * csqrt(x - y * 1j) = conj(csqrt(x + y * 1j)) - // * If x is either +0 or -0 and y is +0, the result is +0 + 0j. - // * If x is any value (including NaN) and y is +infinity, the result - // is +infinity + infinity j. - // * If x is a finite number and y is NaN, the result is NaN + NaN j. - - // * If x -infinity and y is a positive (i.e., greater than 0) finite - // number, the result is NaN + NaN j. - // * If x is +infinity and y is a positive (i.e., greater than 0) - // finite number, the result is +0 + infinity j. - // * If x is -infinity and y is NaN, the result is NaN + infinity j - // (sign of the imaginary component is unspecified). - // * If x is +infinity and y is NaN, the result is +infinity + NaN j. - // * If x is NaN and y is any value, the result is NaN + NaN j. - - using realT = T; - constexpr realT q_nan = std::numeric_limits::quiet_NaN(); - constexpr realT p_inf = std::numeric_limits::infinity(); - constexpr realT zero = realT(0); - - realT x = std::real(z); - realT y = std::imag(z); - - if (std::isinf(y)) { - return {p_inf, y}; - } - else if (std::isnan(x)) { - return {x, q_nan}; - } - else if (std::isinf(x)) { // x is an infinity - // y is either finite, or nan - if (std::signbit(x)) { // x == -inf - return {(std::isfinite(y) ? zero : y), - sycl::copysign(p_inf, y)}; - } - else { - return {p_inf, - (std::isfinite(y) ? sycl::copysign(zero, y) : y)}; - } - } - else { // x is finite - if (std::isfinite(y)) { -#ifdef USE_STD_SQRT_FOR_COMPLEX_TYPES - return std::sqrt(z); -#else - return csqrt_finite(x, y); -#endif - } - else { - return {q_nan, y}; - } - } - } - - int get_normal_scale_float(const float &v) const - { - constexpr int float_significant_bits = 23; - constexpr std::uint32_t exponent_mask = 0xff; - constexpr int exponent_bias = 127; - const int scale = static_cast( - (sycl::bit_cast(v) >> float_significant_bits) & - exponent_mask); - return scale - exponent_bias; - } - - int get_normal_scale_double(const double &v) const - { - constexpr int double_significant_bits = 52; - constexpr std::uint64_t exponent_mask = 0x7ff; - constexpr int exponent_bias = 1023; - const int scale = static_cast( - (sycl::bit_cast(v) >> double_significant_bits) & - exponent_mask); - return scale - exponent_bias; - } - - template int get_normal_scale(const T &v) const - { - static_assert(std::is_same_v || std::is_same_v); - - if constexpr (std::is_same_v) { - return get_normal_scale_float(v); - } - else { - return get_normal_scale_double(v); - } - } - - template - std::complex csqrt_finite_scaled(T const &x, T const &y) const - { - // csqrt(x + y*1j) = - // sqrt((cabs(x, y) + x) / 2) + - // 1j * copysign(sqrt((cabs(x, y) - x) / 2), y) - - using realT = T; - constexpr realT half = realT(0x1.0p-1f); // 1/2 - constexpr realT zero = realT(0); - - const int exp_x = get_normal_scale(x); - const int exp_y = get_normal_scale(y); - - int sc = std::max(exp_x, exp_y) / 2; - const realT xx = sycl::ldexp(x, -sc * 2); - const realT yy = sycl::ldexp(y, -sc * 2); - - if (std::signbit(xx)) { - const realT m = std::hypot(xx, yy); - const realT d = std::sqrt((m - xx) * half); - const realT res_re = (d == zero ? zero : sycl::fabs(yy) / d * half); - const realT res_im = sycl::copysign(d, yy); - return {sycl::ldexp(res_re, sc), sycl::ldexp(res_im, sc)}; - } - else { - const realT m = std::hypot(xx, yy); - const realT d = std::sqrt((m + xx) * half); - const realT res_im = - (d == zero) ? sycl::copysign(zero, yy) : yy * half / d; - return {sycl::ldexp(d, sc), sycl::ldexp(res_im, sc)}; - } - } - - template - std::complex csqrt_finite_unscaled(T const &x, T const &y) const - { - // csqrt(x + y*1j) = - // sqrt((cabs(x, y) + x) / 2) + - // 1j * copysign(sqrt((cabs(x, y) - x) / 2), y) - - using realT = T; - constexpr realT half = realT(0x1.0p-1f); // 1/2 - constexpr realT zero = realT(0); - - if (std::signbit(x)) { - const realT m = std::hypot(x, y); - const realT d = std::sqrt((m - x) * half); - const realT res_re = (d == zero ? zero : sycl::fabs(y) / d * half); - const realT res_im = sycl::copysign(d, y); - return {res_re, res_im}; - } - else { - const realT m = std::hypot(x, y); - const realT d = std::sqrt((m + x) * half); - const realT res_im = - (d == zero) ? sycl::copysign(zero, y) : y * half / d; - return {d, res_im}; - } - } - - template T scaling_threshold() const - { - if constexpr (std::is_same_v) { - return T(0x1.0p+126f); - } - else { - return T(0x1.0p+1022); - } - } - - template - std::complex csqrt_finite(T const &x, T const &y) const - { - return (std::max(std::fabs(x), std::fabs(y)) < - scaling_threshold()) - ? csqrt_finite_unscaled(x, y) - : csqrt_finite_scaled(x, y); - } }; template struct SquareFunctor resT operator()(const argT &in) const { if constexpr (is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using realT = typename argT::value_type; auto z = exprm_ns::complex(in); return z * z; -#else - return in * in; -#endif } else { return in * in; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sycl_complex.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sycl_complex.hpp index 57dc97bb1e..014a0a84d0 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sycl_complex.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sycl_complex.hpp @@ -1,13 +1,10 @@ #pragma once -#ifdef USE_SYCL_FOR_COMPLEX_TYPES + #define SYCL_EXT_ONEAPI_COMPLEX #if __has_include() #include #else #include #endif -#endif -#ifdef USE_SYCL_FOR_COMPLEX_TYPES namespace exprm_ns = sycl::ext::oneapi::experimental; -#endif diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp index 960a441371..c68f8c2910 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp @@ -102,7 +102,7 @@ template struct TanFunctor } const realT tanh_re = sycl::copysign(realT(1), x); const realT tanh_im = sycl::copysign( - realT(0), std::isinf(y) ? y : std::sin(y) * std::cos(y)); + realT(0), std::isinf(y) ? y : sycl::sin(y) * sycl::cos(y)); return resT{tanh_im, -tanh_re}; } /* @@ -118,16 +118,12 @@ template struct TanFunctor return resT{q_nan, q_nan}; } /* ordinary cases */ -#ifdef USE_SYCL_FOR_COMPLEX_TYPES - return exprm_ns::tan(exprm_ns::complex(in)); // std::tan(in); -#else - return std::tan(in); -#endif + return exprm_ns::tan(exprm_ns::complex(in)); // tan(in); } else { static_assert(std::is_floating_point_v || std::is_same_v); - return std::tan(in); + return sycl::tan(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp index 4e0ef989aa..bf4a84fcaf 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp @@ -96,7 +96,7 @@ template struct TanhFunctor } const realT res_re = sycl::copysign(realT(1), x); const realT res_im = sycl::copysign( - realT(0), std::isinf(y) ? y : std::sin(y) * std::cos(y)); + realT(0), std::isinf(y) ? y : sycl::sin(y) * sycl::cos(y)); return resT{res_re, res_im}; } /* @@ -112,17 +112,13 @@ template struct TanhFunctor return resT{q_nan, q_nan}; } /* ordinary cases */ -#ifdef USE_SYCL_FOR_COMPLEX_TYPES return exprm_ns::tanh( - exprm_ns::complex(in)); // std::tanh(in); -#else - return std::tanh(in); -#endif + exprm_ns::complex(in)); // tanh(in); } else { static_assert(std::is_floating_point_v || std::is_same_v); - return std::tanh(in); + return sycl::tanh(in); } } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/true_divide.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/true_divide.hpp index ae38415737..fd2b4a2ddd 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/true_divide.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/true_divide.hpp @@ -64,37 +64,25 @@ struct TrueDivideFunctor if constexpr (tu_ns::is_complex::value && tu_ns::is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using realT1 = typename argT1::value_type; using realT2 = typename argT2::value_type; return exprm_ns::complex(in1) / exprm_ns::complex(in2); -#else - return in1 / in2; -#endif } else if constexpr (tu_ns::is_complex::value && !tu_ns::is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using realT1 = typename argT1::value_type; return exprm_ns::complex(in1) / in2; -#else - return in1 / in2; -#endif } else if constexpr (!tu_ns::is_complex::value && tu_ns::is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using realT2 = typename argT2::value_type; return in1 / exprm_ns::complex(in2); -#else - return in1 / in2; -#endif } else { return in1 / in2; @@ -422,27 +410,19 @@ template struct TrueDivideInplaceFunctor { if constexpr (tu_ns::is_complex::value) { if constexpr (tu_ns::is_complex::value) { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using res_rT = typename resT::value_type; using arg_rT = typename argT::value_type; auto res1 = exprm_ns::complex(res); res1 /= exprm_ns::complex(in); res = res1; -#else - res /= in; -#endif } else { -#ifdef USE_SYCL_FOR_COMPLEX_TYPES using res_rT = typename resT::value_type; auto res1 = exprm_ns::complex(res); res1 /= in; res = res1; -#else - res /= in; -#endif } } else { diff --git a/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp b/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp index 2d411057a1..64a69717eb 100644 --- a/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp +++ b/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp @@ -67,7 +67,7 @@ class WrapIndex void operator()(ssize_t max_item, ssize_t &ind) const { max_item = std::max(max_item, 1); - ind = std::clamp(ind, -max_item, max_item - 1); + ind = sycl::clamp(ind, -max_item, max_item - 1); ind = (ind < 0) ? ind + max_item : ind; return; } @@ -81,7 +81,7 @@ class ClipIndex void operator()(ssize_t max_item, ssize_t &ind) const { max_item = std::max(max_item, 1); - ind = std::clamp(ind, 0, max_item - 1); + ind = sycl::clamp(ind, 0, max_item - 1); return; } }; From 6e226772a9bc5bb2494eb0d867923fec62ee419b Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Sat, 8 Jun 2024 19:07:13 -0500 Subject: [PATCH 2/6] Replace 0 with ssize_t(0) --- .../libtensor/include/kernels/integer_advanced_indexing.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp b/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp index 64a69717eb..ac4cf5c3c0 100644 --- a/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp +++ b/dpctl/tensor/libtensor/include/kernels/integer_advanced_indexing.hpp @@ -81,7 +81,7 @@ class ClipIndex void operator()(ssize_t max_item, ssize_t &ind) const { max_item = std::max(max_item, 1); - ind = sycl::clamp(ind, 0, max_item - 1); + ind = sycl::clamp(ind, ssize_t(0), max_item - 1); return; } }; From 69ce77665c0fcd72568f3ab04015fc39042a7709 Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sun, 9 Jun 2024 03:31:42 +0000 Subject: [PATCH 3/6] Fix pre-commit --- .../libtensor/include/kernels/elementwise_functions/acos.hpp | 3 +-- .../libtensor/include/kernels/elementwise_functions/acosh.hpp | 4 ++-- .../libtensor/include/kernels/elementwise_functions/asinh.hpp | 3 +-- .../libtensor/include/kernels/elementwise_functions/atan.hpp | 3 +-- .../libtensor/include/kernels/elementwise_functions/atanh.hpp | 3 +-- .../libtensor/include/kernels/elementwise_functions/cos.hpp | 3 +-- .../libtensor/include/kernels/elementwise_functions/sin.hpp | 4 ++-- .../libtensor/include/kernels/elementwise_functions/tanh.hpp | 3 +-- 8 files changed, 10 insertions(+), 16 deletions(-) diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp index 0befcd8b74..d2beed632a 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp @@ -116,8 +116,7 @@ template struct AcosFunctor } /* ordinary cases */ - return exprm_ns::acos( - exprm_ns::complex(in)); // acos(in); + return exprm_ns::acos(exprm_ns::complex(in)); // acos(in); } else { static_assert(std::is_floating_point_v || diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp index 1c18d08d35..f187aabb27 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp @@ -120,8 +120,8 @@ template struct AcoshFunctor } else { /* ordinary cases */ - acos_in = exprm_ns::acos( - exprm_ns::complex(in)); // acos(in); + acos_in = + exprm_ns::acos(exprm_ns::complex(in)); // acos(in); } /* Now we calculate acosh(z) */ diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp index 312dd6e43a..f3662776e9 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp @@ -119,8 +119,7 @@ template struct AsinhFunctor } /* ordinary cases */ - return exprm_ns::asinh( - exprm_ns::complex(in)); // asinh(in); + return exprm_ns::asinh(exprm_ns::complex(in)); // asinh(in); } else { static_assert(std::is_floating_point_v || diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp index f7b57c3487..99a6a9a987 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp @@ -126,8 +126,7 @@ template struct AtanFunctor return resT{atanh_im, atanh_re}; } /* ordinary cases */ - return exprm_ns::atan( - exprm_ns::complex(in)); // atan(in); + return exprm_ns::atan(exprm_ns::complex(in)); // atan(in); } else { static_assert(std::is_floating_point_v || diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp index 2fcee9ad14..f8aca408a6 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp @@ -120,8 +120,7 @@ template struct AtanhFunctor return resT{res_re, res_im}; } /* ordinary cases */ - return exprm_ns::atanh( - exprm_ns::complex(in)); // atanh(in); + return exprm_ns::atanh(exprm_ns::complex(in)); // atanh(in); } else { static_assert(std::is_floating_point_v || diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp index 74ba53e40b..308f6b0d44 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp @@ -81,8 +81,7 @@ template struct CosFunctor * real and imaginary parts of input are finite. */ if (in_re_finite && in_im_finite) { - return exprm_ns::cos( - exprm_ns::complex(in)); // cos(in); + return exprm_ns::cos(exprm_ns::complex(in)); // cos(in); } /* diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp index 2bc38bfc12..238981af53 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp @@ -79,8 +79,8 @@ template struct SinFunctor * real and imaginary parts of input are finite. */ if (in_re_finite && in_im_finite) { - resT res = exprm_ns::sin( - exprm_ns::complex(in)); // sin(in); + resT res = + exprm_ns::sin(exprm_ns::complex(in)); // sin(in); if (in_re == realT(0)) { res.real(sycl::copysign(realT(0), in_re)); } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp index bf4a84fcaf..6e638e6186 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp @@ -112,8 +112,7 @@ template struct TanhFunctor return resT{q_nan, q_nan}; } /* ordinary cases */ - return exprm_ns::tanh( - exprm_ns::complex(in)); // tanh(in); + return exprm_ns::tanh(exprm_ns::complex(in)); // tanh(in); } else { static_assert(std::is_floating_point_v || From 163bd3c8819c34f1b885e43b462a218e44a497de Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Sun, 9 Jun 2024 08:55:28 -0500 Subject: [PATCH 4/6] Replaced remaining uses of std::cos, std::exp, etc. --- .../kernels/elementwise_functions/cosh.hpp | 2 +- .../kernels/elementwise_functions/exp.hpp | 4 ++-- .../kernels/elementwise_functions/exp2.hpp | 4 ++-- .../kernels/elementwise_functions/expm1.hpp | 20 +++++++++---------- .../elementwise_functions/logaddexp.hpp | 4 ++-- .../kernels/elementwise_functions/sin.hpp | 4 ++-- .../kernels/elementwise_functions/sinh.hpp | 2 +- .../libtensor/include/utils/math_utils.hpp | 6 +++--- 8 files changed, 23 insertions(+), 23 deletions(-) diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp index 2c3c2643cb..96de1639ca 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp @@ -213,7 +213,7 @@ template struct CoshContigFactory template struct CoshTypeMapFactory { - /*! @brief get typeid for output type of std::cosh(T x) */ + /*! @brief get typeid for output type of sycl::cosh(T x) */ std::enable_if_t::value, int> get() { using rT = typename CoshOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp index fbca696a1a..180fe3a348 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp @@ -104,7 +104,7 @@ template struct ExpFunctor } else { /* x is -inf */ if (std::isfinite(y)) { - realT exp_x = std::exp(x); + realT exp_x = sycl::exp(x); return resT{exp_x * sycl::cos(y), exp_x * sycl::sin(y)}; } else { @@ -182,7 +182,7 @@ template struct ExpContigFactory template struct ExpTypeMapFactory { - /*! @brief get typeid for output type of std::exp(T x) */ + /*! @brief get typeid for output type of sycl::exp(T x) */ std::enable_if_t::value, int> get() { using rT = typename ExpOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp index 5bc652bd0d..2609088b75 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp @@ -106,7 +106,7 @@ template struct Exp2Functor } else { /* x is -inf */ if (std::isfinite(y)) { - realT exp_x = std::exp(x); + realT exp_x = sycl::exp(x); return resT{exp_x * sycl::cos(y), exp_x * sycl::sin(y)}; } else { @@ -184,7 +184,7 @@ template struct Exp2ContigFactory template struct Exp2TypeMapFactory { - /*! @brief get typeid for output type of std::exp2(T x) */ + /*! @brief get typeid for output type of sycl::exp2(T x) */ std::enable_if_t::value, int> get() { using rT = typename Exp2OutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/expm1.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/expm1.hpp index 114a0cb417..1002ad1f1e 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/expm1.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/expm1.hpp @@ -84,8 +84,8 @@ template struct Expm1Functor return in; } else { - return (resT{sycl::copysign(x, std::cos(y)), - sycl::copysign(x, std::sin(y))}); + return (resT{sycl::copysign(x, sycl::cos(y)), + sycl::copysign(x, sycl::sin(y))}); } } else { @@ -97,7 +97,7 @@ template struct Expm1Functor } else { return resT{realT(-1), - sycl::copysign(realT(0), std::sin(y))}; + sycl::copysign(realT(0), sycl::sin(y))}; } } } @@ -113,13 +113,13 @@ template struct Expm1Functor } // x, y finite numbers - const realT cosY_val = std::cos(y); - const realT sinY_val = (y == 0) ? y : std::sin(y); - const realT sinhalfY_val = (y == 0) ? y : std::sin(y / 2); + const realT cosY_val = sycl::cos(y); + const realT sinY_val = (y == 0) ? y : sycl::sin(y); + const realT sinhalfY_val = (y == 0) ? y : sycl::sin(y / 2); const realT res_re = - std::expm1(x) * cosY_val - 2 * sinhalfY_val * sinhalfY_val; - realT res_im = std::exp(x) * sinY_val; + sycl::expm1(x) * cosY_val - 2 * sinhalfY_val * sinhalfY_val; + realT res_im = sycl::exp(x) * sinY_val; return resT{res_re, res_im}; } else { @@ -129,7 +129,7 @@ template struct Expm1Functor if (in == 0) { return in; } - return std::expm1(in); + return sycl::expm1(in); } } }; @@ -197,7 +197,7 @@ template struct Expm1ContigFactory template struct Expm1TypeMapFactory { - /*! @brief get typeid for output type of std::expm1(T x) */ + /*! @brief get typeid for output type of sycl::expm1(T x) */ std::enable_if_t::value, int> get() { using rT = typename Expm1OutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logaddexp.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logaddexp.hpp index 4276893690..842b30248a 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logaddexp.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logaddexp.hpp @@ -91,8 +91,8 @@ template struct LogAddExpFunctor private: template T impl_finite(T const &in) const { - return (in > 0) ? (in + std::log1p(std::exp(-in))) - : std::log1p(std::exp(in)); + return (in > 0) ? (in + sycl::log1p(sycl::exp(-in))) + : sycl::log1p(sycl::exp(in)); } }; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp index 238981af53..d8425cd5a0 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp @@ -158,7 +158,7 @@ template struct SinFunctor } const realT sinh_re = x * sycl::cos(y); const realT sinh_im = - std::numeric_limits::infinity() * std::sin(y); + std::numeric_limits::infinity() * sycl::sin(y); return resT{sinh_im, -sinh_re}; } @@ -246,7 +246,7 @@ template struct SinContigFactory template struct SinTypeMapFactory { - /*! @brief get typeid for output type of std::sin(T x) */ + /*! @brief get typeid for output type of sycl::sin(T x) */ std::enable_if_t::value, int> get() { using rT = typename SinOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp index 64dbde7eef..c5ed20ad1b 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp @@ -215,7 +215,7 @@ template struct SinhContigFactory template struct SinhTypeMapFactory { - /*! @brief get typeid for output type of std::sinh(T x) */ + /*! @brief get typeid for output type of sycl::sinh(T x) */ std::enable_if_t::value, int> get() { using rT = typename SinhOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/utils/math_utils.hpp b/dpctl/tensor/libtensor/include/utils/math_utils.hpp index 7038d46193..5f2ce8b01e 100644 --- a/dpctl/tensor/libtensor/include/utils/math_utils.hpp +++ b/dpctl/tensor/libtensor/include/utils/math_utils.hpp @@ -118,16 +118,16 @@ template T min_complex(const T &x1, const T &x2) template T logaddexp(T x, T y) { if (x == y) { // handle signed infinities - const T log2 = std::log(T(2)); + const T log2 = sycl::log(T(2)); return x + log2; } else { const T tmp = x - y; if (tmp > 0) { - return x + std::log1p(std::exp(-tmp)); + return x + sycl::log1p(sycl::exp(-tmp)); } else if (tmp <= 0) { - return y + std::log1p(std::exp(tmp)); + return y + sycl::log1p(sycl::exp(tmp)); } else { return std::numeric_limits::quiet_NaN(); From 0650dc891be6ee6b43548d1162737e21ba15a27f Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sun, 9 Jun 2024 21:38:55 +0000 Subject: [PATCH 5/6] Revert `sycl::log1p` to `std::log1p` in `math_utils.hpp` until compiler crash on CUDA is resolved --- dpctl/tensor/libtensor/include/utils/math_utils.hpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/dpctl/tensor/libtensor/include/utils/math_utils.hpp b/dpctl/tensor/libtensor/include/utils/math_utils.hpp index 5f2ce8b01e..97e2fee377 100644 --- a/dpctl/tensor/libtensor/include/utils/math_utils.hpp +++ b/dpctl/tensor/libtensor/include/utils/math_utils.hpp @@ -122,12 +122,14 @@ template T logaddexp(T x, T y) return x + log2; } else { + // FIXME: switch to `sycl::log1p` when + // compiler segfault in CUDA build is fixed const T tmp = x - y; if (tmp > 0) { - return x + sycl::log1p(sycl::exp(-tmp)); + return x + std::log1p(sycl::exp(-tmp)); } else if (tmp <= 0) { - return y + sycl::log1p(sycl::exp(tmp)); + return y + std::log1p(sycl::exp(tmp)); } else { return std::numeric_limits::quiet_NaN(); From 6ea1120edce21942dd50733c904a4576f06177cc Mon Sep 17 00:00:00 2001 From: Nikita Grigorian Date: Sun, 9 Jun 2024 21:40:21 +0000 Subject: [PATCH 6/6] Replace remaining `std` namespace transcendental functions A few stray instances of `signbit`, `atan`, `log`, and `exp` as well as `atan2` and `hypot` --- .../include/kernels/elementwise_functions/acos.hpp | 2 +- .../include/kernels/elementwise_functions/acosh.hpp | 2 +- .../include/kernels/elementwise_functions/asin.hpp | 4 ++-- .../include/kernels/elementwise_functions/asinh.hpp | 2 +- .../include/kernels/elementwise_functions/atan.hpp | 2 +- .../include/kernels/elementwise_functions/atan2.hpp | 6 +++--- .../include/kernels/elementwise_functions/atanh.hpp | 6 +++--- .../include/kernels/elementwise_functions/exp.hpp | 2 +- .../include/kernels/elementwise_functions/exp2.hpp | 4 ++-- .../include/kernels/elementwise_functions/hypot.hpp | 4 ++-- .../include/kernels/elementwise_functions/log.hpp | 2 +- .../include/kernels/elementwise_functions/log10.hpp | 4 ++-- .../include/kernels/elementwise_functions/log1p.hpp | 12 ++++++------ .../include/kernels/elementwise_functions/log2.hpp | 2 +- .../kernels/elementwise_functions/logaddexp.hpp | 2 +- .../include/kernels/elementwise_functions/tan.hpp | 2 +- .../include/kernels/elementwise_functions/tanh.hpp | 2 +- 17 files changed, 30 insertions(+), 30 deletions(-) diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp index d2beed632a..5b8204f8b1 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp @@ -188,7 +188,7 @@ template struct AcosContigFactory template 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::value, int> get() { using rT = typename AcosOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp index f187aabb27..6ecfe02765 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp @@ -215,7 +215,7 @@ template struct AcoshContigFactory template 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::value, int> get() { using rT = typename AcoshOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp index 49f0513bf0..e98693cdfc 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp @@ -136,7 +136,7 @@ template struct AsinFunctor } /* ordinary cases */ return exprm_ns::asin( - exprm_ns::complex(in)); // std::asin(in); + exprm_ns::complex(in)); // sycl::asin(in); } else { static_assert(std::is_floating_point_v || @@ -208,7 +208,7 @@ template struct AsinContigFactory template 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::value, int> get() { using rT = typename AsinOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp index f3662776e9..9850ab6a3c 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp @@ -191,7 +191,7 @@ template struct AsinhContigFactory template 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::value, int> get() { using rT = typename AsinhOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp index 99a6a9a987..0e78c01903 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp @@ -198,7 +198,7 @@ template struct AtanContigFactory template 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::value, int> get() { using rT = typename AtanOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan2.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan2.hpp index 3b587a69f4..f5cacb9178 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan2.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan2.hpp @@ -56,12 +56,12 @@ template 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); } }; @@ -145,7 +145,7 @@ template struct Atan2ContigFactory template 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::value, int> get() { using rT = typename Atan2OutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp index f8aca408a6..c9130619a5 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp @@ -76,7 +76,7 @@ template struct AtanhFunctor 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 res_re = sycl::copysign(realT(0), x); const realT res_im = sycl::copysign(pi_half, y); @@ -113,7 +113,7 @@ template struct AtanhFunctor realT(1) / std::numeric_limits::epsilon(); if (sycl::fabs(x) > RECIP_EPSILON || sycl::fabs(y) > RECIP_EPSILON) { - const realT pi_half = std::atan(realT(1)) * 2; + const realT pi_half = sycl::atan(realT(1)) * 2; const realT res_re = realT(0); const realT res_im = sycl::copysign(pi_half, y); @@ -192,7 +192,7 @@ template struct AtanhContigFactory template struct AtanhTypeMapFactory { - /*! @brief get typeid for output type of std::atanh(T x) */ + /*! @brief get typeid for output type of sycl::atanh(T x) */ std::enable_if_t::value, int> get() { using rT = typename AtanhOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp index 180fe3a348..cd451dda5a 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp @@ -90,7 +90,7 @@ template struct ExpFunctor } } else { - if (!std::signbit(x)) { /* x is +inf */ + if (!sycl::signbit(x)) { /* x is +inf */ if (y == realT(0)) { return resT{x, y}; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp index 2609088b75..dbe32b47d1 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp @@ -68,7 +68,7 @@ template struct Exp2Functor if constexpr (is_complex::value) { using realT = typename argT::value_type; - const argT tmp = in * std::log(realT(2)); + const argT tmp = in * sycl::log(realT(2)); constexpr realT q_nan = std::numeric_limits::quiet_NaN(); @@ -92,7 +92,7 @@ template struct Exp2Functor } } else { - if (!std::signbit(x)) { /* x is +inf */ + if (!sycl::signbit(x)) { /* x is +inf */ if (y == realT(0)) { return resT{x, y}; } diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/hypot.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/hypot.hpp index 7f79d2e6eb..53817373a0 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/hypot.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/hypot.hpp @@ -58,7 +58,7 @@ template struct HypotFunctor resT operator()(const argT1 &in1, const argT2 &in2) const { - return std::hypot(in1, in2); + return sycl::hypot(in1, in2); } template @@ -160,7 +160,7 @@ template struct HypotContigFactory template struct HypotTypeMapFactory { - /*! @brief get typeid for output type of std::hypot(T1 x, T2 y) */ + /*! @brief get typeid for output type of sycl::hypot(T1 x, T2 y) */ std::enable_if_t::value, int> get() { using rT = typename HypotOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp index fa2bc20b2f..5fddfd9b64 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp @@ -139,7 +139,7 @@ template struct LogContigFactory template struct LogTypeMapFactory { - /*! @brief get typeid for output type of std::log(T x) */ + /*! @brief get typeid for output type of sycl::log(T x) */ std::enable_if_t::value, int> get() { using rT = typename LogOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp index a7ce78a757..d062da6c6c 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp @@ -73,7 +73,7 @@ template struct Log10Functor using realT = typename argT::value_type; // return (log(in) / log(realT{10})); return exprm_ns::log(exprm_ns::complex(in)) / - std::log(realT{10}); + sycl::log(realT{10}); } else { return sycl::log10(in); @@ -158,7 +158,7 @@ template struct Log10ContigFactory template struct Log10TypeMapFactory { - /*! @brief get typeid for output type of std::log10(T x) */ + /*! @brief get typeid for output type of sycl::log10(T x) */ std::enable_if_t::value, int> get() { using rT = typename Log10OutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log1p.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log1p.hpp index 4b524504fc..9cc0927cca 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log1p.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log1p.hpp @@ -79,23 +79,23 @@ template struct Log1pFunctor const realT y = std::imag(in); // imaginary part of result - const realT res_im = std::atan2(y, x + 1); + const realT res_im = sycl::atan2(y, x + 1); if (std::max(sycl::fabs(x), sycl::fabs(y)) < realT{.1}) { const realT v = x * (2 + x) + y * y; - return resT{std::log1p(v) / 2, res_im}; + return resT{sycl::log1p(v) / 2, res_im}; } else { // when not close to zero, // prevent overflow - const realT m = std::hypot(x + 1, y); - return resT{std::log(m), res_im}; + const realT m = sycl::hypot(x + 1, y); + return resT{sycl::log(m), res_im}; } } else { static_assert(std::is_floating_point_v || std::is_same_v); - return std::log1p(in); + return sycl::log1p(in); } } }; @@ -163,7 +163,7 @@ template struct Log1pContigFactory template struct Log1pTypeMapFactory { - /*! @brief get typeid for output type of std::log1p(T x) */ + /*! @brief get typeid for output type of sycl::log1p(T x) */ std::enable_if_t::value, int> get() { using rT = typename Log1pOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp index b98d93f6eb..dac784c7a4 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp @@ -159,7 +159,7 @@ template struct Log2ContigFactory template struct Log2TypeMapFactory { - /*! @brief get typeid for output type of std::log2(T x) */ + /*! @brief get typeid for output type of sycl::log2(T x) */ std::enable_if_t::value, int> get() { using rT = typename Log2OutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logaddexp.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logaddexp.hpp index 842b30248a..d94724edc6 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logaddexp.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/logaddexp.hpp @@ -177,7 +177,7 @@ template struct LogAddExpContigFactory template struct LogAddExpTypeMapFactory { - /*! @brief get typeid for output type of std::logaddexp(T1 x, T2 y) */ + /*! @brief get typeid for output type of logaddexp(T1 x, T2 y) */ std::enable_if_t::value, int> get() { using rT = typename LogAddExpOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp index c68f8c2910..03bb5baba1 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp @@ -190,7 +190,7 @@ template struct TanContigFactory template struct TanTypeMapFactory { - /*! @brief get typeid for output type of std::tan(T x) */ + /*! @brief get typeid for output type of sycl::tan(T x) */ std::enable_if_t::value, int> get() { using rT = typename TanOutputType::value_type; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp index 6e638e6186..db275bc5ea 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp @@ -184,7 +184,7 @@ template struct TanhContigFactory template struct TanhTypeMapFactory { - /*! @brief get typeid for output type of std::tanh(T x) */ + /*! @brief get typeid for output type of sycl::tanh(T x) */ std::enable_if_t::value, int> get() { using rT = typename TanhOutputType::value_type;