Skip to content

Use sycl experimental complex namespace for real and imag #2062

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

Closed
wants to merge 2 commits into from
Closed

Use sycl experimental complex namespace for real and imag #2062

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
0a447ad
Update implementations of `real` and `imag`
ndgrigorian Apr 25, 2025
7564d1c
Use `is_complex_v` in real and imag
ndgrigorian Apr 27, 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
16 changes: 10 additions & 6 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/imag.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,7 @@
#include <sycl/sycl.hpp>
#include <type_traits>

#include "sycl_complex.hpp"
#include "vec_size_util.hpp"

#include "kernels/dpctl_tensor_types.hpp"
Expand All @@ -53,14 +54,15 @@ using dpctl::tensor::ssize_t;
namespace td_ns = dpctl::tensor::type_dispatch;

using dpctl::tensor::type_utils::is_complex;
using dpctl::tensor::type_utils::is_complex_v;

template <typename argT, typename resT> struct ImagFunctor
{

// is function constant for given argT
using is_constant = typename std::false_type;
using is_constant =
typename std::is_same<is_complex<argT>, std::false_type>;
// constant value, if constant
// constexpr resT constant_value = resT{};
static constexpr resT constant_value = resT{0};
// is function defined for sycl::vec
using supports_vec = typename std::false_type;
// do both argTy and resTy support sugroup store/load operation
Expand All @@ -69,12 +71,14 @@ template <typename argT, typename resT> struct ImagFunctor

resT operator()(const argT &in) const
{
if constexpr (is_complex<argT>::value) {
return std::imag(in);
if constexpr (is_complex_v<argT>) {
using realT = typename argT::value_type;
using sycl_complexT = typename exprm_ns::complex<realT>;
return exprm_ns::imag(sycl_complexT(in));
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It might be unclear from the PR description what was the motivation to switch to experimental namespace and why other place where std::imag used are not impacted.. Would it make sense to add clarity on that?

Copy link
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The initial motivation was old technical debt to add the constant value for imag, since we may as well use it for real-valued types.

In terms of why add in the experimental namespace change, complex header is not supported on AMD and this has been an issue with recent AMD builds. I've pushed another branch which removes (almost) every use of complex in kernel code aside from conversion to SYCL complex type. I want to see if this resolves issues with AMD builds on OS compiler.

Copy link
Collaborator Author

@ndgrigorian ndgrigorian Apr 26, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In general, I think it would be best to slowly transition to full use of the sycl complex extension. We also have benchmarking now and can see if this PR brings regressions once we have some unary elementwise benchmarks in place.

But it could be done in a separate PR, if you feel it makes more sense that way.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I see now, thank you. I'm totally with that. Just it was a bit unclear from the PR description.

}
else {
static_assert(std::is_same_v<resT, argT>);
return resT{0};
return constant_value;
}
}
};
Expand Down
8 changes: 6 additions & 2 deletions dpctl/tensor/libtensor/include/kernels/elementwise_functions/real.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,7 @@
#include <sycl/sycl.hpp>
#include <type_traits>

#include "sycl_complex.hpp"
#include "vec_size_util.hpp"

#include "kernels/dpctl_tensor_types.hpp"
Expand All @@ -53,6 +54,7 @@ using dpctl::tensor::ssize_t;
namespace td_ns = dpctl::tensor::type_dispatch;

using dpctl::tensor::type_utils::is_complex;
using dpctl::tensor::type_utils::is_complex_v;

template <typename argT, typename resT> struct RealFunctor
{
Expand All @@ -69,8 +71,10 @@ template <typename argT, typename resT> struct RealFunctor

resT operator()(const argT &in) const
{
if constexpr (is_complex<argT>::value) {
return std::real(in);
if constexpr (is_complex_v<argT>) {
using realT = typename argT::value_type;
using sycl_complexT = typename exprm_ns::complex<realT>;
return exprm_ns::real(sycl_complexT(in));
}
else {
static_assert(std::is_same_v<resT, argT>);
Expand Down
Loading