From 0f7689088a7deb340407f60411f91437eacf12d0 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk Date: Thu, 16 Nov 2023 16:51:02 -0600 Subject: [PATCH] Allow for change of name and location of sycl_complex.hpp Introduced private header to load SYCL's experimental complex header from the right location. The header and implementations respond to USE_SYCL_FOR_COMPLEX_TYPES preprocessor variable. If set, sycl::ext::oneapi::experimental namespace functions are to be used. Otherwise std:: namespace functions will be used instead for complex types. USE_SYCL_FOR_COMPLEX_TYPES is being set in tensor/CMakeLists.txt --- dpctl/tensor/CMakeLists.txt | 4 +-- .../kernels/elementwise_functions/abs.hpp | 5 ++-- .../kernels/elementwise_functions/acos.hpp | 17 ++++++++++-- .../kernels/elementwise_functions/acosh.hpp | 13 +++++++-- .../kernels/elementwise_functions/add.hpp | 15 +++++++++-- .../kernels/elementwise_functions/asin.hpp | 26 +++++++++++++++--- .../kernels/elementwise_functions/asinh.hpp | 13 +++++++-- .../kernels/elementwise_functions/atan.hpp | 7 +++-- .../kernels/elementwise_functions/atanh.hpp | 7 +++-- .../kernels/elementwise_functions/conj.hpp | 7 +++-- .../kernels/elementwise_functions/cos.hpp | 7 +++-- .../kernels/elementwise_functions/cosh.hpp | 7 +++-- .../kernels/elementwise_functions/equal.hpp | 7 +++-- .../kernels/elementwise_functions/exp.hpp | 7 +++-- .../kernels/elementwise_functions/exp2.hpp | 7 +++-- .../kernels/elementwise_functions/log.hpp | 8 ++++-- .../kernels/elementwise_functions/log10.hpp | 8 ++++-- .../kernels/elementwise_functions/log2.hpp | 8 ++++-- .../elementwise_functions/multiply.hpp | 8 ++++-- .../kernels/elementwise_functions/pow.hpp | 11 ++++++-- .../kernels/elementwise_functions/sign.hpp | 7 +++-- .../kernels/elementwise_functions/sin.hpp | 7 +++-- .../kernels/elementwise_functions/sinh.hpp | 7 +++-- .../kernels/elementwise_functions/sqrt.hpp | 17 ++++++------ .../kernels/elementwise_functions/square.hpp | 7 +++-- .../elementwise_functions/sycl_complex.hpp | 13 +++++++++ .../kernels/elementwise_functions/tan.hpp | 9 ++++--- .../kernels/elementwise_functions/tanh.hpp | 11 +++++--- .../elementwise_functions/true_divide.hpp | 27 ++++++++++++++++--- 29 files changed, 228 insertions(+), 69 deletions(-) create mode 100644 dpctl/tensor/libtensor/include/kernels/elementwise_functions/sycl_complex.hpp diff --git a/dpctl/tensor/CMakeLists.txt b/dpctl/tensor/CMakeLists.txt index f2454a9fdc..d6346bfd50 100644 --- a/dpctl/tensor/CMakeLists.txt +++ b/dpctl/tensor/CMakeLists.txt @@ -188,9 +188,9 @@ foreach(_src_fn ${_no_fast_math_sources}) ) endforeach() if (UNIX) - set(_compiler_definitions "USE_STD_ABS_FOR_COMPLEX_TYPES;USE_STD_SQRT_FOR_COMPLEX_TYPES;SYCL_EXT_ONEAPI_COMPLEX") + set(_compiler_definitions "USE_SYCL_FOR_COMPLEX_TYPES") else() - set(_compiler_definitions "SYCL_EXT_ONEAPI_COMPLEX") + set(_compiler_definitions "USE_SYCL_FOR_COMPLEX_TYPES") endif() foreach(_src_fn ${_elementwise_sources}) diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp index ab321ad356..911452931e 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/abs.hpp @@ -28,11 +28,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -50,7 +50,6 @@ namespace abs namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -121,7 +120,7 @@ template struct AbsFunctor return q_nan; } else { -#ifdef USE_STD_ABS_FOR_COMPLEX_TYPES +#ifdef USE_SYCL_FOR_COMPLEX_TYPES return exprm_ns::abs(exprm_ns::complex(z)); #else return std::hypot(std::real(z), std::imag(z)); diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp index 23a87b9d44..c4742e66dc 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp @@ -26,11 +26,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -48,7 +48,6 @@ namespace acos namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -105,6 +104,7 @@ template struct AcosFunctor constexpr realT r_eps = realT(1) / std::numeric_limits::epsilon(); if (std::abs(x) > r_eps || std::abs(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)); @@ -115,11 +115,24 @@ template struct AcosFunctor 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 = std::abs(wy); + + realT ry = wx + std::log(realT(2)); + return resT{rx, (std::signbit(y)) ? ry : -ry}; +#endif } /* 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 } 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 56730a411c..b736d5b658 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp @@ -26,11 +26,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -48,7 +48,6 @@ namespace acosh namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -112,18 +111,28 @@ template struct AcoshFunctor * For large x or y including acos(+-Inf + I*+-Inf) */ if (std::abs(x) > r_eps || std::abs(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 = std::abs(wy); realT ry = wx + std::log(realT(2)); acos_in = resT{rx, (std::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 } /* Now we calculate acosh(z) */ diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp index 0ed1710833..1297847831 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp @@ -26,10 +26,10 @@ #pragma once #include #include -#include #include #include +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" #include "utils/type_utils.hpp" @@ -50,7 +50,6 @@ namespace add namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; namespace tu_ns = dpctl::tensor::type_utils; -namespace exprm_ns = sycl::ext::oneapi::experimental; template struct AddFunctor { @@ -65,24 +64,36 @@ 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/asin.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp index 035480c437..9da5077665 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp @@ -26,11 +26,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -48,7 +48,6 @@ namespace asin namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -119,26 +118,45 @@ template struct AsinFunctor constexpr realT r_eps = realT(1) / std::numeric_limits::epsilon(); if (std::abs(x) > r_eps || std::abs(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)) { - auto log_z = exprm_ns::log(z); + const auto log_z = exprm_ns::log(z); wx = log_z.real() + std::log(realT(2)); wy = log_z.imag(); } else { - auto log_mz = exprm_ns::log(-z); + const auto log_mz = exprm_ns::log(-z); wx = log_mz.real() + std::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 = std::copysign(wx, x); const realT asinh_im = std::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 || diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp index 523ca4f01f..ceab296b91 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp @@ -26,11 +26,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -48,7 +48,6 @@ namespace asinh namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -108,20 +107,30 @@ template struct AsinhFunctor realT(1) / std::numeric_limits::epsilon(); if (std::abs(x) > r_eps || std::abs(y) > r_eps) { +#ifdef USE_SYCL_FOR_COMPLEX_TYPES using sycl_complexT = exprm_ns::complex; sycl_complexT log_in = (std::signbit(x)) ? exprm_ns::log(sycl_complexT(-in)) : exprm_ns::log(sycl_complexT(in)); realT wx = log_in.real() + std::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 = std::copysign(wx, x); const realT res_im = std::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 } 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 df8bba538b..a9af6d829e 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp @@ -27,11 +27,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -49,7 +49,6 @@ namespace atan namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -128,8 +127,12 @@ template struct AtanFunctor 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 } 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 d6a4b06ac3..3be6abb742 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp @@ -27,11 +27,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -49,7 +49,6 @@ namespace atanh namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -121,8 +120,12 @@ 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 } else { static_assert(std::is_floating_point_v || diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp index 6977e3a747..a3c607aa49 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp @@ -28,11 +28,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -50,7 +50,6 @@ namespace conj namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -70,9 +69,13 @@ 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 } 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 bdc1acc1fe..c8cd8ef18c 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp @@ -26,11 +26,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -48,7 +48,6 @@ namespace cos namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -83,8 +82,12 @@ 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 } /* diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp index 7093d2a2a3..f03d438cbe 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp @@ -26,11 +26,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -48,7 +48,6 @@ namespace cosh namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -83,8 +82,12 @@ 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 } /* diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/equal.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/equal.hpp index 6d68861396..9f866aa580 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/equal.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/equal.hpp @@ -26,10 +26,10 @@ #pragma once #include #include -#include #include #include +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" #include "utils/type_utils.hpp" @@ -49,7 +49,6 @@ namespace equal namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; namespace tu_ns = dpctl::tensor::type_utils; -namespace exprm_ns = sycl::ext::oneapi::experimental; template struct EqualFunctor { @@ -70,8 +69,12 @@ 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 { return (in1 == in2); diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp index 453eb05c52..add15c3523 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp.hpp @@ -26,11 +26,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -48,7 +48,6 @@ namespace exp namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -75,8 +74,12 @@ 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 } else { return resT{q_nan, q_nan}; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp index b6b2f32e83..a22411414b 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/exp2.hpp @@ -27,11 +27,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -49,7 +49,6 @@ namespace exp2 namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -78,7 +77,11 @@ 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}; diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp index ff37d87157..3668551473 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log.hpp @@ -24,13 +24,14 @@ #pragma once #include +#include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -48,7 +49,6 @@ namespace log namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -68,8 +68,12 @@ 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 } else { return std::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 88dabcaabe..5997aded5f 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log10.hpp @@ -25,13 +25,14 @@ #pragma once #include +#include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -49,7 +50,6 @@ namespace log10 namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; using dpctl::tensor::type_utils::vec_cast; @@ -72,9 +72,13 @@ 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 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); diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp index 57d7dcaf31..211a5fdb6c 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/log2.hpp @@ -25,13 +25,14 @@ #pragma once #include +#include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -49,7 +50,6 @@ namespace log2 namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; using dpctl::tensor::type_utils::vec_cast; @@ -72,9 +72,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}); return exprm_ns::log(exprm_ns::complex(in)) / std::log(realT{2}); +#else + return std::log(in) / std::log(realT{2}); +#endif } else { return std::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 612ad78360..47549b5a74 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/multiply.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/multiply.hpp @@ -24,12 +24,13 @@ //===---------------------------------------------------------------------===// #pragma once +#include #include #include -#include #include #include +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" #include "utils/type_utils.hpp" @@ -50,7 +51,6 @@ namespace multiply namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; namespace tu_ns = dpctl::tensor::type_utils; -namespace exprm_ns = sycl::ext::oneapi::experimental; template struct MultiplyFunctor { @@ -65,11 +65,15 @@ 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 95e8442903..9068e67f10 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/pow.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/pow.hpp @@ -27,10 +27,10 @@ #include #include #include -#include #include #include +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" #include "utils/type_utils.hpp" @@ -51,7 +51,6 @@ namespace pow namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; namespace tu_ns = dpctl::tensor::type_utils; -namespace exprm_ns = sycl::ext::oneapi::experimental; template struct PowFunctor { @@ -88,11 +87,15 @@ 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); @@ -365,11 +368,15 @@ 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); diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sign.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sign.hpp index 162db394de..521e935c16 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sign.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sign.hpp @@ -27,11 +27,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -49,7 +49,6 @@ namespace sign namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; using dpctl::tensor::type_utils::vec_cast; @@ -81,8 +80,12 @@ template struct SignFunctor return resT(0); } else { +#ifdef USE_SYCL_FOR_COMPLEX_TYPES auto z = exprm_ns::complex(in); return (z / exprm_ns::abs(z)); +#else + return in / std::abs(in); +#endif } } else { diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp index e1e9e79c57..97768fc8e9 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp @@ -26,11 +26,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -48,7 +48,6 @@ namespace sin namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -81,8 +80,12 @@ template struct SinFunctor * real and imaginary parts of input are finite. */ if (in_re_finite && in_im_finite) { +#ifdef USE_SYCL_FOR_COMPLEX_TYPES return exprm_ns::sin( exprm_ns::complex(in)); // std::sin(in); +#else + return std::sin(in); +#endif } /* diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp index b11c7402d0..fe7dae533b 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp @@ -26,11 +26,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -48,7 +48,6 @@ namespace sinh namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -81,7 +80,11 @@ 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. diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sqrt.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sqrt.hpp index b638e4a55f..70ace858b5 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sqrt.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sqrt.hpp @@ -29,11 +29,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -51,7 +51,6 @@ namespace sqrt namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -71,15 +70,17 @@ template struct SqrtFunctor resT operator()(const argT &in) const { if constexpr (is_complex::value) { - // #ifdef _WINDOWS - // return csqrt(in); - // #else - // return std::sqrt(in); - // #endif +#ifdef USE_SYCL_FOR_COMPLEX_TYPES using realT = typename argT::value_type; - // return csqrt(in); return exprm_ns::sqrt(exprm_ns::complex(in)); +#else +#ifdef _WINDOWS + return std::sqrt(in); +#else + return std::sqrt(in); +#endif +#endif } else { return std::sqrt(in); diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/square.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/square.hpp index 2c37ce87d9..c888c4eefa 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/square.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/square.hpp @@ -27,11 +27,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -49,7 +49,6 @@ namespace square namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace exprm_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; using dpctl::tensor::type_utils::vec_cast; @@ -71,11 +70,15 @@ 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 new file mode 100644 index 0000000000..57dc97bb1e --- /dev/null +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/sycl_complex.hpp @@ -0,0 +1,13 @@ +#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 1f97b59054..2c648edd7b 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp @@ -27,11 +27,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -49,7 +49,6 @@ namespace tan namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace cmplx_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -120,7 +119,11 @@ template struct TanFunctor return resT{q_nan, q_nan}; } /* ordinary cases */ - return cmplx_ns::tan(cmplx_ns::complex(in)); // std::tan(in); +#ifdef USE_SYCL_FOR_COMPLEX_TYPES + return exprm_ns::tan(exprm_ns::complex(in)); // std::tan(in); +#else + return std::tan(in); +#endif } else { static_assert(std::is_floating_point_v || diff --git a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp index 453ce17b54..84ba7a9a3f 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp @@ -28,11 +28,11 @@ #include #include #include -#include #include #include #include "kernels/elementwise_functions/common.hpp" +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" @@ -50,7 +50,6 @@ namespace tanh namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; -namespace cmplx_ns = sycl::ext::oneapi::experimental; using dpctl::tensor::type_utils::is_complex; @@ -114,8 +113,12 @@ template struct TanhFunctor return resT{q_nan, q_nan}; } /* ordinary cases */ - return cmplx_ns::tanh( - cmplx_ns::complex(in)); // std::tanh(in); +#ifdef USE_SYCL_FOR_COMPLEX_TYPES + return exprm_ns::tanh( + exprm_ns::complex(in)); // std::tanh(in); +#else + return std::tanh(in); +#endif } else { static_assert(std::is_floating_point_v || 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 6620d2e3c1..c72634d106 100644 --- a/dpctl/tensor/libtensor/include/kernels/elementwise_functions/true_divide.hpp +++ b/dpctl/tensor/libtensor/include/kernels/elementwise_functions/true_divide.hpp @@ -26,10 +26,10 @@ #pragma once #include #include -#include #include #include +#include "sycl_complex.hpp" #include "utils/offset_utils.hpp" #include "utils/type_dispatch.hpp" #include "utils/type_utils.hpp" @@ -50,7 +50,6 @@ namespace true_divide namespace py = pybind11; namespace td_ns = dpctl::tensor::type_dispatch; namespace tu_ns = dpctl::tensor::type_utils; -namespace exprm_ns = sycl::ext::oneapi::experimental; template struct TrueDivideFunctor @@ -66,25 +65,37 @@ 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; @@ -409,18 +420,28 @@ template struct TrueDivideInplaceFunctor void operator()(resT &res, const argT &in) { if constexpr (tu_ns::is_complex::value) { - using res_rT = typename resT::value_type; 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 {