Skip to content

Use experimental complex extension for all complex arithmetic #2069

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
ndgrigorian wants to merge 9 commits into
base: master
Choose a base branch
from
+747 −389
Draft

Use experimental complex extension for all complex arithmetic #2069

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
ae75273
Use sycl complex extension throughout element-wise and utils
ndgrigorian Apr 26, 2025
88e64c1
Update binary functions multiply and subtract to use experimental SYC…
ndgrigorian Apr 29, 2025
01f22d3
Use experimental SYCL complex in dot product
ndgrigorian Apr 29, 2025
34168bd
Use experimental namespace in sequential dot product
ndgrigorian Apr 30, 2025
0fd49ae
Use experimental complex namespace in gemm
ndgrigorian Apr 30, 2025
e3be74e
Use specialized functor for multiplying or adding complex inputs
ndgrigorian Apr 30, 2025
6fba5f2
Use custom plus operator in gemm and dot product tree reduction kernels
ndgrigorian Apr 30, 2025
da15d0e
Refactor operators when dispatching to tree reductions
ndgrigorian Apr 30, 2025
52bb73e
Refactor sycl_complex indirect include
ndgrigorian Apr 30, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
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
  • Loading branch information
@ndgrigorian
ndgrigorian committed May 1, 2025
commit 52bb73e7704d3b73e401b482ee429e4c013a4c54
6 changes: 4 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/acos.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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;

Expand All @@ -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);
Expand Down
6 changes: 4 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/acosh.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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;

Expand Down Expand Up @@ -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);
Expand Down
27 changes: 22 additions & 5 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/add.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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
{
Expand All @@ -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;
Expand Down Expand Up @@ -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,
Expand Down
6 changes: 4 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/angle.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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;

Expand All @@ -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);
}
};

Expand Down
6 changes: 4 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/asin.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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;

Expand Down Expand Up @@ -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);
Expand Down
6 changes: 4 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/asinh.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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;

Expand All @@ -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);
Expand Down
6 changes: 4 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/atan.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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;
Expand Down Expand Up @@ -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);
Expand Down
6 changes: 4 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/atanh.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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;

Expand All @@ -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);
Expand Down
9 changes: 6 additions & 3 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/cabs_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,7 @@
#include <complex>
#include <limits>

#include "sycl_complex.hpp"
#include "utils/sycl_complex.hpp"

namespace dpctl
{
Expand All @@ -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):
Expand All @@ -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);

Expand Down
6 changes: 4 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/conj.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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;

Expand All @@ -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>)
Expand Down
6 changes: 4 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/cos.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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;

Expand All @@ -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);
Expand Down
6 changes: 4 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/cosh.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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;

Expand All @@ -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);
Expand Down
Loading
Loading