Refactor sycl_complex indirect include

* Move sycl_complex.hpp to utils
* No longer use exprm_ns defined by header, define on per-file basis
* Include alias to type sycl_complex_t<T> under sycl_utils namespace
* Use identical include macro where inclusion of sycl_complex would be impossible

Author: ndgrigorian

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

@@ -29,7 +29,7 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "sycl_complex.hpp"
+#include "utils/sycl_complex.hpp"
 #include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
@@ -49,7 +49,9 @@ namespace acos
 {
 
 using dpctl::tensor::ssize_t;
+namespace su_ns = dpctl::tensor::sycl_utils;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace exprm_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -72,7 +74,7 @@ template <typename argT, typename resT> struct AcosFunctor
             using realT = typename argT::value_type;
 
             constexpr realT q_nan = std::numeric_limits<realT>::quiet_NaN();
-            using sycl_complexT = exprm_ns::complex<realT>;
+            using sycl_complexT = su_ns::sycl_complex_t<realT>;
             sycl_complexT z = sycl_complexT(in);
             const realT x = exprm_ns::real(z);
             const realT y = exprm_ns::imag(z);

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

@@ -29,7 +29,7 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "sycl_complex.hpp"
+#include "utils/sycl_complex.hpp"
 #include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
@@ -49,7 +49,9 @@ namespace acosh
 {
 
 using dpctl::tensor::ssize_t;
+namespace su_ns = dpctl::tensor::sycl_utils;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace exprm_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -77,7 +79,7 @@ template <typename argT, typename resT> struct AcoshFunctor
              * where the sign is chosen so Re(acosh(in)) >= 0.
              * So, we first calculate acos(in) and then acosh(in).
              */
-            using sycl_complexT = exprm_ns::complex<realT>;
+            using sycl_complexT = su_ns::sycl_complex_t<realT>;
             sycl_complexT z = sycl_complexT(in);
             const realT x = exprm_ns::real(z);
             const realT y = exprm_ns::imag(z);

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

@@ -29,7 +29,7 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "sycl_complex.hpp"
+#include "utils/sycl_complex.hpp"
 #include "vec_size_util.hpp"
 
 #include "utils/offset_utils.hpp"
@@ -50,8 +50,10 @@ namespace add
 {
 
 using dpctl::tensor::ssize_t;
+namespace su_ns = dpctl::tensor::sycl_utils;
 namespace td_ns = dpctl::tensor::type_dispatch;
 namespace tu_ns = dpctl::tensor::type_utils;
+namespace exprm_ns = sycl::ext::oneapi::experimental;
 
 template <typename argT1, typename argT2, typename resT> struct AddFunctor
 {
@@ -69,21 +71,22 @@ template <typename argT1, typename argT2, typename resT> struct AddFunctor
             using rT1 = typename argT1::value_type;
             using rT2 = typename argT2::value_type;
 
-            return exprm_ns::complex<rT1>(in1) + exprm_ns::complex<rT2>(in2);
+            return su_ns::sycl_complex_t<rT1>(in1) +
+                   su_ns::sycl_complex_t<rT2>(in2);
         }
         else if constexpr (tu_ns::is_complex<argT1>::value &&
                            !tu_ns::is_complex<argT2>::value)
         {
             using rT1 = typename argT1::value_type;
 
-            return exprm_ns::complex<rT1>(in1) + in2;
+            return su_ns::sycl_complex_t<rT1>(in1) + in2;
         }
         else if constexpr (!tu_ns::is_complex<argT1>::value &&
                            tu_ns::is_complex<argT2>::value)
         {
             using rT2 = typename argT2::value_type;
 
-            return in1 + exprm_ns::complex<rT2>(in2);
+            return in1 + su_ns::sycl_complex_t<rT2>(in2);
         }
         else {
             return in1 + in2;
@@ -460,7 +463,21 @@ template <typename argT, typename resT> struct AddInplaceFunctor
     using supports_vec = std::negation<
         std::disjunction<tu_ns::is_complex<argT>, tu_ns::is_complex<resT>>>;
 
-    void operator()(resT &res, const argT &in) { res += in; }
+    void operator()(resT &res, const argT &in)
+    {
+        if constexpr (tu_ns::is_complex_v<resT> && tu_ns::is_complex_v<argT>) {
+            using rT1 = typename resT::value_type;
+            using rT2 = typename argT::value_type;
+
+            auto tmp = su_ns::sycl_complex_t<rT1>(res);
+            tmp += su_ns::sycl_complex_t<rT2>(in);
+
+            res = resT(tmp);
+        }
+        else {
+            res += in;
+        }
+    }
 
     template <int vec_sz>
     void operator()(sycl::vec<resT, vec_sz> &res,

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

@@ -30,7 +30,7 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "sycl_complex.hpp"
+#include "utils/sycl_complex.hpp"
 #include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
@@ -50,7 +50,9 @@ namespace angle
 {
 
 using dpctl::tensor::ssize_t;
+namespace su_ns = dpctl::tensor::sycl_utils;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace exprm_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -71,7 +73,7 @@ template <typename argT, typename resT> struct AngleFunctor
     {
         using rT = typename argT::value_type;
 
-        return exprm_ns::arg(exprm_ns::complex<rT>(in)); // arg(in);
+        return exprm_ns::arg(su_ns::sycl_complex_t<rT>(in)); // arg(in);
     }
 };
 

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

@@ -29,7 +29,7 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "sycl_complex.hpp"
+#include "utils/sycl_complex.hpp"
 #include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
@@ -49,7 +49,9 @@ namespace asin
 {
 
 using dpctl::tensor::ssize_t;
+namespace su_ns = dpctl::tensor::sycl_utils;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace exprm_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -80,7 +82,7 @@ template <typename argT, typename resT> struct AsinFunctor
              * y = imag(I * conj(in)) = real(in)
              * and then return {imag(w), real(w)} which is asin(in)
              */
-            using sycl_complexT = exprm_ns::complex<realT>;
+            using sycl_complexT = su_ns::sycl_complex_t<realT>;
             sycl_complexT z = sycl_complexT(in);
             const realT x = exprm_ns::imag(z);
             const realT y = exprm_ns::real(z);

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

@@ -29,7 +29,7 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "sycl_complex.hpp"
+#include "utils/sycl_complex.hpp"
 #include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
@@ -49,7 +49,9 @@ namespace asinh
 {
 
 using dpctl::tensor::ssize_t;
+namespace su_ns = dpctl::tensor::sycl_utils;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace exprm_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -72,7 +74,7 @@ template <typename argT, typename resT> struct AsinhFunctor
             using realT = typename argT::value_type;
 
             constexpr realT q_nan = std::numeric_limits<realT>::quiet_NaN();
-            using sycl_complexT = exprm_ns::complex<realT>;
+            using sycl_complexT = su_ns::sycl_complex_t<realT>;
             sycl_complexT z = sycl_complexT(in);
             const realT x = exprm_ns::real(z);
             const realT y = exprm_ns::imag(z);

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

@@ -30,7 +30,7 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "sycl_complex.hpp"
+#include "utils/sycl_complex.hpp"
 #include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
@@ -50,7 +50,9 @@ namespace atan
 {
 
 using dpctl::tensor::ssize_t;
+namespace su_ns = dpctl::tensor::sycl_utils;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace exprm_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::kernels::vec_size_utils::ContigHyperparameterSetDefault;
 using dpctl::tensor::kernels::vec_size_utils::UnaryContigHyperparameterSetEntry;
@@ -83,7 +85,7 @@ template <typename argT, typename resT> struct AtanFunctor
              * y = imag(I * conj(in)) = real(in)
              * and then return {imag(w), real(w)} which is atan(in)
              */
-            using sycl_complexT = exprm_ns::complex<realT>;
+            using sycl_complexT = su_ns::sycl_complex_t<realT>;
             sycl_complexT z = sycl_complexT(in);
             const realT x = exprm_ns::imag(z);
             const realT y = exprm_ns::real(z);

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

@@ -30,7 +30,7 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "sycl_complex.hpp"
+#include "utils/sycl_complex.hpp"
 #include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
@@ -50,7 +50,9 @@ namespace atanh
 {
 
 using dpctl::tensor::ssize_t;
+namespace su_ns = dpctl::tensor::sycl_utils;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace exprm_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -73,7 +75,7 @@ template <typename argT, typename resT> struct AtanhFunctor
             using realT = typename argT::value_type;
             constexpr realT q_nan = std::numeric_limits<realT>::quiet_NaN();
 
-            using sycl_complexT = exprm_ns::complex<realT>;
+            using sycl_complexT = su_ns::sycl_complex_t<realT>;
             sycl_complexT z = sycl_complexT(in);
             const realT x = exprm_ns::real(z);
             const realT y = exprm_ns::imag(z);

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

@@ -27,7 +27,7 @@
 #include <complex>
 #include <limits>
 
-#include "sycl_complex.hpp"
+#include "utils/sycl_complex.hpp"
 
 namespace dpctl
 {
@@ -38,6 +38,9 @@ namespace kernels
 namespace detail
 {
 
+namespace su_ns = dpctl::tensor::sycl_utils;
+namespace exprm_ns = sycl::ext::oneapi::experimental;
+
 template <typename realT> realT cabs(std::complex<realT> const &z)
 {
     // Special values for cabs( x + y * 1j):
@@ -51,8 +54,8 @@ template <typename realT> realT cabs(std::complex<realT> const &z)
     //   * If x is a finite number and y is NaN, the result is NaN.
     //   * If x is NaN and y is NaN, the result is NaN.
 
-    using sycl_complexT = exprm_ns::complex<realT>;
-    sycl_complexT _z = exprm_ns::complex<realT>(z);
+    using sycl_complexT = su_ns::sycl_complex_t<realT>;
+    sycl_complexT _z = su_ns::sycl_complex_t<realT>(z);
     const realT x = exprm_ns::real(_z);
     const realT y = exprm_ns::imag(_z);
 

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

@@ -31,7 +31,7 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "sycl_complex.hpp"
+#include "utils/sycl_complex.hpp"
 #include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
@@ -51,7 +51,9 @@ namespace conj
 {
 
 using dpctl::tensor::ssize_t;
+namespace su_ns = dpctl::tensor::sycl_utils;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace exprm_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -73,7 +75,7 @@ template <typename argT, typename resT> struct ConjFunctor
         if constexpr (is_complex<argT>::value) {
             using rT = typename argT::value_type;
 
-            return exprm_ns::conj(exprm_ns::complex<rT>(in)); // conj(in);
+            return exprm_ns::conj(su_ns::sycl_complex_t<rT>(in)); // conj(in);
         }
         else {
             if constexpr (!std::is_same_v<argT, bool>)

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

@@ -29,7 +29,7 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "sycl_complex.hpp"
+#include "utils/sycl_complex.hpp"
 #include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
@@ -49,7 +49,9 @@ namespace cos
 {
 
 using dpctl::tensor::ssize_t;
+namespace su_ns = dpctl::tensor::sycl_utils;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace exprm_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -72,7 +74,7 @@ template <typename argT, typename resT> struct CosFunctor
             using realT = typename argT::value_type;
 
             constexpr realT q_nan = std::numeric_limits<realT>::quiet_NaN();
-            using sycl_complexT = exprm_ns::complex<realT>;
+            using sycl_complexT = su_ns::sycl_complex_t<realT>;
             sycl_complexT z = sycl_complexT(in);
             const realT z_re = exprm_ns::real(z);
             const realT z_im = exprm_ns::imag(z);

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

@@ -29,7 +29,7 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "sycl_complex.hpp"
+#include "utils/sycl_complex.hpp"
 #include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
@@ -49,7 +49,9 @@ namespace cosh
 {
 
 using dpctl::tensor::ssize_t;
+namespace su_ns = dpctl::tensor::sycl_utils;
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace exprm_ns = sycl::ext::oneapi::experimental;
 
 using dpctl::tensor::type_utils::is_complex;
 
@@ -73,7 +75,7 @@ template <typename argT, typename resT> struct CoshFunctor
 
             constexpr realT q_nan = std::numeric_limits<realT>::quiet_NaN();
 
-            using sycl_complexT = exprm_ns::complex<realT>;
+            using sycl_complexT = su_ns::sycl_complex_t<realT>;
             sycl_complexT z = sycl_complexT(in);
             const realT x = exprm_ns::real(z);
             const realT y = exprm_ns::imag(z);