Use sycl::ext::oneapi::experimental for complex trig/trigh and inverses

Use sycl_complex extension to implement complex-valued trigonometric,
hyperbolic functions and their inverses.

This works around use of double precision functions/literals in implementations
of these functions in MSVC headers, causing failures to offload on Iris Xe for
single precision input citing lack of fp64 support by the hardware.

Author: oleksandr-pavlyk

dpctl/tensor/CMakeLists.txt

@@ -64,7 +64,7 @@ set_source_files_properties(
 if (UNIX)
   set_source_files_properties(
     ${CMAKE_CURRENT_SOURCE_DIR}/libtensor/source/elementwise_functions.cpp
-    PROPERTIES COMPILE_DEFINITIONS "USE_STD_ABS_FOR_COMPLEX_TYPES;USE_STD_SQRT_FOR_COMPLEX_TYPES")
+    PROPERTIES COMPILE_DEFINITIONS "USE_STD_ABS_FOR_COMPLEX_TYPES;USE_STD_SQRT_FOR_COMPLEX_TYPES;SYCL_EXT_ONEAPI_COMPLEX")
 endif()
 target_compile_options(${python_module_name} PRIVATE -fno-sycl-id-queries-fit-in-int)
 target_link_options(${python_module_name} PRIVATE -fsycl-device-code-split=per_kernel)

dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp

@@ -27,6 +27,7 @@
 #include <cmath>
 #include <cstddef>
 #include <cstdint>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
 #include <type_traits>
 
 #include "kernels/elementwise_functions/common.hpp"
@@ -47,6 +48,7 @@ namespace acos
 
 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,7 +116,8 @@ template <typename argT, typename resT> struct AcosFunctor
             }
 
             /* ordinary cases */
-            return std::acos(in);
+            return cmplx_ns::acos(
+                cmplx_ns::complex<realT>(in)); // std::acos(in);
         }
         else {
             static_assert(std::is_floating_point_v<argT> ||

dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp

@@ -27,6 +27,7 @@
 #include <cmath>
 #include <cstddef>
 #include <cstdint>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
 #include <type_traits>
 
 #include "kernels/elementwise_functions/common.hpp"
@@ -47,6 +48,7 @@ namespace acosh
 
 namespace py = pybind11;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace cmplx_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -118,7 +120,8 @@ template <typename argT, typename resT> struct AcoshFunctor
             }
             else {
                 /* ordinary cases */
-                acos_in = std::acos(in);
+                acos_in = cmplx_ns::acos(
+                    cmplx_ns::complex<realT>(in)); // std::acos(in);
             }
 
             /* Now we calculate acosh(z) */

dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp

@@ -27,6 +27,7 @@
 #include <cmath>
 #include <cstddef>
 #include <cstdint>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
 #include <type_traits>
 
 #include "kernels/elementwise_functions/common.hpp"
@@ -47,6 +48,7 @@ namespace asin
 
 namespace py = pybind11;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace cmplx_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -134,7 +136,8 @@ template <typename argT, typename resT> struct AsinFunctor
                 return resT{asinh_im, asinh_re};
             }
             /* ordinary cases */
-            return std::asin(in);
+            return cmplx_ns::asin(
+                cmplx_ns::complex<realT>(in)); // std::asin(in);
         }
         else {
             static_assert(std::is_floating_point_v<argT> ||

dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp

@@ -27,6 +27,7 @@
 #include <cmath>
 #include <cstddef>
 #include <cstdint>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
 #include <type_traits>
 
 #include "kernels/elementwise_functions/common.hpp"
@@ -47,6 +48,7 @@ namespace asinh
 
 namespace py = pybind11;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace cmplx_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -115,7 +117,8 @@ template <typename argT, typename resT> struct AsinhFunctor
             }
 
             /* ordinary cases */
-            return std::asinh(in);
+            return cmplx_ns::asinh(
+                cmplx_ns::complex<realT>(in)); // std::asinh(in);
         }
         else {
             static_assert(std::is_floating_point_v<argT> ||

dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp

@@ -28,6 +28,7 @@
 #include <complex>
 #include <cstddef>
 #include <cstdint>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
 #include <type_traits>
 
 #include "kernels/elementwise_functions/common.hpp"
@@ -48,6 +49,7 @@ namespace atan
 
 namespace py = pybind11;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace cmplx_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -126,7 +128,8 @@ template <typename argT, typename resT> struct AtanFunctor
                 return resT{atanh_im, atanh_re};
             }
             /* ordinary cases */
-            return std::atan(in);
+            return cmplx_ns::atan(
+                cmplx_ns::complex<realT>(in)); // std::atan(in);
         }
         else {
             static_assert(std::is_floating_point_v<argT> ||

dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp

@@ -28,6 +28,7 @@
 #include <complex>
 #include <cstddef>
 #include <cstdint>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
 #include <type_traits>
 
 #include "kernels/elementwise_functions/common.hpp"
@@ -48,6 +49,7 @@ namespace atanh
 
 namespace py = pybind11;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace cmplx_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -119,7 +121,8 @@ template <typename argT, typename resT> struct AtanhFunctor
                 return resT{res_re, res_im};
             }
             /* ordinary cases */
-            return std::atanh(in);
+            return cmplx_ns::atanh(
+                cmplx_ns::complex<realT>(in)); // std::atanh(in);
         }
         else {
             static_assert(std::is_floating_point_v<argT> ||

dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp

@@ -27,6 +27,7 @@
 #include <cmath>
 #include <cstddef>
 #include <cstdint>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
 #include <type_traits>
 
 #include "kernels/elementwise_functions/common.hpp"
@@ -47,6 +48,7 @@ namespace cos
 
 namespace py = pybind11;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace cmplx_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -81,7 +83,8 @@ template <typename argT, typename resT> struct CosFunctor
              * real and imaginary parts of input are finite.
              */
             if (in_re_finite && in_im_finite) {
-                return std::cos(in);
+                return cmplx_ns::cos(
+                    cmplx_ns::complex<realT>(in)); // std::cos(in);
             }
 
             /*

dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp

@@ -27,6 +27,7 @@
 #include <cmath>
 #include <cstddef>
 #include <cstdint>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
 #include <type_traits>
 
 #include "kernels/elementwise_functions/common.hpp"
@@ -47,6 +48,7 @@ namespace cosh
 
 namespace py = pybind11;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace cmplx_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -81,7 +83,8 @@ template <typename argT, typename resT> struct CoshFunctor
              * real and imaginary parts of input are finite.
              */
             if (xfinite && yfinite) {
-                return std::cosh(in);
+                return cmplx_ns::cosh(
+                    cmplx_ns::complex<realT>(in)); // std::cosh(in);
             }
 
             /*

dpctl/tensor/libtensor/include/kernels/elementwise_functions/sin.hpp

@@ -23,10 +23,11 @@
 //===---------------------------------------------------------------------===//
 
 #pragma once
-#include <CL/sycl.hpp>
 #include <cmath>
 #include <cstddef>
 #include <cstdint>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
+#include <sycl/sycl.hpp>
 #include <type_traits>
 
 #include "kernels/elementwise_functions/common.hpp"
@@ -47,6 +48,7 @@ namespace sin
 
 namespace py = pybind11;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace cmplx_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -79,7 +81,8 @@ template <typename argT, typename resT> struct SinFunctor
              * real and imaginary parts of input are finite.
              */
             if (in_re_finite && in_im_finite) {
-                return std::sin(in);
+                return cmplx_ns::sin(
+                    cmplx_ns::complex<realT>(in)); // std::sin(in);
             }
 
             /*

dpctl/tensor/libtensor/include/kernels/elementwise_functions/sinh.hpp

@@ -27,6 +27,7 @@
 #include <cmath>
 #include <cstddef>
 #include <cstdint>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
 #include <type_traits>
 
 #include "kernels/elementwise_functions/common.hpp"
@@ -47,6 +48,7 @@ namespace sinh
 
 namespace py = pybind11;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace cmplx_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 

dpctl/tensor/libtensor/include/kernels/elementwise_functions/tan.hpp

@@ -28,6 +28,7 @@
 #include <complex>
 #include <cstddef>
 #include <cstdint>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
 #include <type_traits>
 
 #include "kernels/elementwise_functions/common.hpp"
@@ -48,6 +49,7 @@ 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;
 
@@ -118,7 +120,7 @@ template <typename argT, typename resT> struct TanFunctor
                 return resT{q_nan, q_nan};
             }
             /* ordinary cases */
-            return std::tan(in);
+            return cmplx_ns::tan(cmplx_ns::complex<realT>(in)); // std::tan(in);
         }
         else {
             static_assert(std::is_floating_point_v<argT> ||

dpctl/tensor/libtensor/include/kernels/elementwise_functions/tanh.hpp

@@ -29,6 +29,7 @@
 #include <complex>
 #include <cstddef>
 #include <cstdint>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
 #include <type_traits>
 
 #include "kernels/elementwise_functions/common.hpp"
@@ -49,6 +50,7 @@ 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;
 
@@ -112,7 +114,8 @@ template <typename argT, typename resT> struct TanhFunctor
                 return resT{q_nan, q_nan};
             }
             /* ordinary cases */
-            return std::tanh(in);
+            return cmplx_ns::tanh(
+                cmplx_ns::complex<realT>(in)); // std::tanh(in);
         }
         else {
             static_assert(std::is_floating_point_v<argT> ||