From 8184818a61e12c0f656a12ea45f55819b2a72a31 Mon Sep 17 00:00:00 2001
From: Hugh Bird <hugh.bird@codeplay.com>
Date: Thu, 30 May 2024 16:32:07 +0100
Subject: [PATCH 1/3] [DFT] Correct overload resolution for OOP COMPLEX vs IP
 REAL_REAL

* OOP COMPLEX and IP REAL_REAL overload resolution is problematic
* Correct with SFINAE
wq
---
 include/oneapi/mkl/dft/backward.hpp          | 14 ++++++---
 include/oneapi/mkl/dft/detail/types_impl.hpp | 32 ++++++++++++++++++++
 include/oneapi/mkl/dft/forward.hpp           | 16 ++++++----
 3 files changed, 52 insertions(+), 10 deletions(-)

diff --git a/include/oneapi/mkl/dft/backward.hpp b/include/oneapi/mkl/dft/backward.hpp
index b63dabb28..5eafe48c5 100644
--- a/include/oneapi/mkl/dft/backward.hpp
+++ b/include/oneapi/mkl/dft/backward.hpp
@@ -44,7 +44,8 @@ void compute_backward(descriptor_type &desc, sycl::buffer<data_type, 1> &inout)
 }
 
 //In-place transform, using config_param::COMPLEX_STORAGE=config_value::REAL_REAL data format
-template <typename descriptor_type, typename data_type>
+template <typename descriptor_type, typename data_type,
+          std::enable_if_t<detail::valid_ip_realreal_impl<descriptor_type, data_type>, bool> = true>
 void compute_backward(descriptor_type &desc, sycl::buffer<data_type, 1> &inout_re,
                       sycl::buffer<data_type, 1> &inout_im) {
     static_assert(detail::valid_compute_arg<descriptor_type, data_type>::value,
@@ -59,7 +60,9 @@ void compute_backward(descriptor_type &desc, sycl::buffer<data_type, 1> &inout_r
 }
 
 //Out-of-place transform
-template <typename descriptor_type, typename input_type, typename output_type>
+template <typename descriptor_type, typename input_type, typename output_type,
+          std::enable_if_t<detail::valid_oop_iotypes<descriptor_type, input_type, output_type>,
+                           bool> = true>
 void compute_backward(descriptor_type &desc, sycl::buffer<input_type, 1> &in,
                       sycl::buffer<output_type, 1> &out) {
     static_assert(detail::valid_compute_arg<descriptor_type, input_type>::value,
@@ -114,7 +117,8 @@ sycl::event compute_backward(descriptor_type &desc, data_type *inout,
 }
 
 //In-place transform, using config_param::COMPLEX_STORAGE=config_value::REAL_REAL data format
-template <typename descriptor_type, typename data_type>
+template <typename descriptor_type, typename data_type,
+          std::enable_if_t<detail::valid_ip_realreal_impl<descriptor_type, data_type>, bool> = true>
 sycl::event compute_backward(descriptor_type &desc, data_type *inout_re, data_type *inout_im,
                              const std::vector<sycl::event> &dependencies = {}) {
     static_assert(detail::valid_compute_arg<descriptor_type, data_type>::value,
@@ -127,7 +131,9 @@ sycl::event compute_backward(descriptor_type &desc, data_type *inout_re, data_ty
 }
 
 //Out-of-place transform
-template <typename descriptor_type, typename input_type, typename output_type>
+template <typename descriptor_type, typename input_type, typename output_type,
+          std::enable_if_t<detail::valid_oop_iotypes<descriptor_type, input_type, output_type>,
+                           bool> = true>
 sycl::event compute_backward(descriptor_type &desc, input_type *in, output_type *out,
                              const std::vector<sycl::event> &dependencies = {}) {
     static_assert(detail::valid_compute_arg<descriptor_type, input_type>::value,
diff --git a/include/oneapi/mkl/dft/detail/types_impl.hpp b/include/oneapi/mkl/dft/detail/types_impl.hpp
index e41a71469..75486530a 100644
--- a/include/oneapi/mkl/dft/detail/types_impl.hpp
+++ b/include/oneapi/mkl/dft/detail/types_impl.hpp
@@ -87,6 +87,20 @@ struct descriptor_info<descriptor<precision::DOUBLE, domain::COMPLEX>> {
     using backward_type = std::complex<double>;
 };
 
+// Get the scalar type associated with a descriptor.
+template <class descriptor_t>
+using descriptor_scalar_t = typename descriptor_info<descriptor_t>::scalar_type;
+
+template <typename T>
+constexpr bool is_complex_dft = false;
+template <precision Prec>
+constexpr bool is_complex_dft<descriptor<Prec, domain::COMPLEX>> = true;
+
+template <typename T>
+constexpr bool is_complex = false;
+template <typename T>
+constexpr bool is_complex<std::complex<T>> = true;
+
 template <typename T, typename... Ts>
 using is_one_of = typename std::bool_constant<(std::is_same_v<T, Ts> || ...)>;
 
@@ -97,6 +111,24 @@ using valid_compute_arg = typename std::bool_constant<
     (std::is_same_v<typename detail::descriptor_info<descriptor_type>::scalar_type, double> &&
      is_one_of<T, double, sycl::double2, sycl::double4, std::complex<double>>::value)>;
 
+// For out-of-place complex-complex DFTs, are the input and output types correct? For SFINAE.
+template <class descriptor_t, typename input_t, typename output_t>
+constexpr bool valid_oop_iotypes = []() {
+    if constexpr (is_complex_dft<descriptor_t>) {
+        // Both input and output types must be complex, otherwise select real-real inplace overload.
+        return is_complex<input_t> && is_complex<output_t>;
+    }
+    else {
+        // I/O can be real or complex - no issues resolving overload with real-real inplace.
+        return valid_compute_arg<descriptor_t, input_t>::value &&
+               valid_compute_arg<descriptor_t, output_t>::value;
+    }
+}();
+
+template <class descriptor_t, typename data_t>
+constexpr bool valid_ip_realreal_impl =
+    is_complex_dft<descriptor_t>&& std::is_same_v<descriptor_scalar_t<descriptor_t>, data_t>;
+
 // compute the range of a reinterpreted buffer
 template <typename In, typename Out>
 std::size_t reinterpret_range(std::size_t size) {
diff --git a/include/oneapi/mkl/dft/forward.hpp b/include/oneapi/mkl/dft/forward.hpp
index 3fcd60b1c..353a5a89f 100644
--- a/include/oneapi/mkl/dft/forward.hpp
+++ b/include/oneapi/mkl/dft/forward.hpp
@@ -45,7 +45,8 @@ void compute_forward(descriptor_type &desc, sycl::buffer<data_type, 1> &inout) {
 }
 
 //In-place transform, using config_param::COMPLEX_STORAGE=config_value::REAL_REAL data format
-template <typename descriptor_type, typename data_type>
+template <typename descriptor_type, typename data_type,
+          std::enable_if_t<detail::valid_ip_realreal_impl<descriptor_type, data_type>, bool> = true>
 void compute_forward(descriptor_type &desc, sycl::buffer<data_type, 1> &inout_re,
                      sycl::buffer<data_type, 1> &inout_im) {
     static_assert(detail::valid_compute_arg<descriptor_type, data_type>::value,
@@ -60,7 +61,9 @@ void compute_forward(descriptor_type &desc, sycl::buffer<data_type, 1> &inout_re
 }
 
 //Out-of-place transform
-template <typename descriptor_type, typename input_type, typename output_type>
+template <typename descriptor_type, typename input_type, typename output_type,
+          std::enable_if_t<detail::valid_oop_iotypes<descriptor_type, input_type, output_type>,
+                           bool> = true>
 void compute_forward(descriptor_type &desc, sycl::buffer<input_type, 1> &in,
                      sycl::buffer<output_type, 1> &out) {
     static_assert(detail::valid_compute_arg<descriptor_type, input_type>::value,
@@ -114,26 +117,27 @@ sycl::event compute_forward(descriptor_type &desc, data_type *inout,
 }
 
 //In-place transform, using config_param::COMPLEX_STORAGE=config_value::REAL_REAL data format
-template <typename descriptor_type, typename data_type>
+template <typename descriptor_type, typename data_type,
+          std::enable_if_t<detail::valid_ip_realreal_impl<descriptor_type, data_type>, bool> = true>
 sycl::event compute_forward(descriptor_type &desc, data_type *inout_re, data_type *inout_im,
                             const std::vector<sycl::event> &dependencies = {}) {
     static_assert(detail::valid_compute_arg<descriptor_type, data_type>::value,
                   "unexpected type for data_type");
-
     using scalar_type = typename detail::descriptor_info<descriptor_type>::scalar_type;
     return get_commit(desc)->forward_ip_rr(desc, reinterpret_cast<scalar_type *>(inout_re),
                                            reinterpret_cast<scalar_type *>(inout_im), dependencies);
 }
 
 //Out-of-place transform
-template <typename descriptor_type, typename input_type, typename output_type>
+template <typename descriptor_type, typename input_type, typename output_type,
+          std::enable_if_t<detail::valid_oop_iotypes<descriptor_type, input_type, output_type>,
+                           bool> = true>
 sycl::event compute_forward(descriptor_type &desc, input_type *in, output_type *out,
                             const std::vector<sycl::event> &dependencies = {}) {
     static_assert(detail::valid_compute_arg<descriptor_type, input_type>::value,
                   "unexpected type for input_type");
     static_assert(detail::valid_compute_arg<descriptor_type, output_type>::value,
                   "unexpected type for output_type");
-
     using fwd_type = typename detail::descriptor_info<descriptor_type>::forward_type;
     using bwd_type = typename detail::descriptor_info<descriptor_type>::backward_type;
     return get_commit(desc)->forward_op_cc(desc, reinterpret_cast<fwd_type *>(in),

From 5ad57d36a85d7b2ebbcf8e0a136127c404095e27 Mon Sep 17 00:00:00 2001
From: Hugh Bird <hugh.bird@codeplay.com>
Date: Fri, 31 May 2024 14:01:13 +0100
Subject: [PATCH 2/3] Use descriptor_scalar_t elsewhere

---
 include/oneapi/mkl/dft/detail/types_impl.hpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/include/oneapi/mkl/dft/detail/types_impl.hpp b/include/oneapi/mkl/dft/detail/types_impl.hpp
index 75486530a..131e27c4d 100644
--- a/include/oneapi/mkl/dft/detail/types_impl.hpp
+++ b/include/oneapi/mkl/dft/detail/types_impl.hpp
@@ -106,9 +106,9 @@ using is_one_of = typename std::bool_constant<(std::is_same_v<T, Ts> || ...)>;
 
 template <typename descriptor_type, typename T>
 using valid_compute_arg = typename std::bool_constant<
-    (std::is_same_v<typename detail::descriptor_info<descriptor_type>::scalar_type, float> &&
+    (std::is_same_v<descriptor_scalar_t<descriptor_type>, float> &&
      is_one_of<T, float, sycl::float2, sycl::float4, std::complex<float>>::value) ||
-    (std::is_same_v<typename detail::descriptor_info<descriptor_type>::scalar_type, double> &&
+    (std::is_same_v<descriptor_scalar_t<descriptor_type>, double> &&
      is_one_of<T, double, sycl::double2, sycl::double4, std::complex<double>>::value)>;
 
 // For out-of-place complex-complex DFTs, are the input and output types correct? For SFINAE.

From 516c7fa16a444a9ed400cfbb2b118cad970831f3 Mon Sep 17 00:00:00 2001
From: Hugh Bird <hugh.bird@codeplay.com>
Date: Thu, 6 Jun 2024 11:43:41 +0100
Subject: [PATCH 3/3] Reenable OOP complex-complex version for float input

---
 include/oneapi/mkl/dft/backward.hpp          |  8 ++------
 include/oneapi/mkl/dft/detail/types_impl.hpp | 14 --------------
 include/oneapi/mkl/dft/forward.hpp           |  8 ++------
 3 files changed, 4 insertions(+), 26 deletions(-)

diff --git a/include/oneapi/mkl/dft/backward.hpp b/include/oneapi/mkl/dft/backward.hpp
index 5eafe48c5..3cd03e13b 100644
--- a/include/oneapi/mkl/dft/backward.hpp
+++ b/include/oneapi/mkl/dft/backward.hpp
@@ -60,9 +60,7 @@ void compute_backward(descriptor_type &desc, sycl::buffer<data_type, 1> &inout_r
 }
 
 //Out-of-place transform
-template <typename descriptor_type, typename input_type, typename output_type,
-          std::enable_if_t<detail::valid_oop_iotypes<descriptor_type, input_type, output_type>,
-                           bool> = true>
+template <typename descriptor_type, typename input_type, typename output_type>
 void compute_backward(descriptor_type &desc, sycl::buffer<input_type, 1> &in,
                       sycl::buffer<output_type, 1> &out) {
     static_assert(detail::valid_compute_arg<descriptor_type, input_type>::value,
@@ -131,9 +129,7 @@ sycl::event compute_backward(descriptor_type &desc, data_type *inout_re, data_ty
 }
 
 //Out-of-place transform
-template <typename descriptor_type, typename input_type, typename output_type,
-          std::enable_if_t<detail::valid_oop_iotypes<descriptor_type, input_type, output_type>,
-                           bool> = true>
+template <typename descriptor_type, typename input_type, typename output_type>
 sycl::event compute_backward(descriptor_type &desc, input_type *in, output_type *out,
                              const std::vector<sycl::event> &dependencies = {}) {
     static_assert(detail::valid_compute_arg<descriptor_type, input_type>::value,
diff --git a/include/oneapi/mkl/dft/detail/types_impl.hpp b/include/oneapi/mkl/dft/detail/types_impl.hpp
index 131e27c4d..6fd07694b 100644
--- a/include/oneapi/mkl/dft/detail/types_impl.hpp
+++ b/include/oneapi/mkl/dft/detail/types_impl.hpp
@@ -111,20 +111,6 @@ using valid_compute_arg = typename std::bool_constant<
     (std::is_same_v<descriptor_scalar_t<descriptor_type>, double> &&
      is_one_of<T, double, sycl::double2, sycl::double4, std::complex<double>>::value)>;
 
-// For out-of-place complex-complex DFTs, are the input and output types correct? For SFINAE.
-template <class descriptor_t, typename input_t, typename output_t>
-constexpr bool valid_oop_iotypes = []() {
-    if constexpr (is_complex_dft<descriptor_t>) {
-        // Both input and output types must be complex, otherwise select real-real inplace overload.
-        return is_complex<input_t> && is_complex<output_t>;
-    }
-    else {
-        // I/O can be real or complex - no issues resolving overload with real-real inplace.
-        return valid_compute_arg<descriptor_t, input_t>::value &&
-               valid_compute_arg<descriptor_t, output_t>::value;
-    }
-}();
-
 template <class descriptor_t, typename data_t>
 constexpr bool valid_ip_realreal_impl =
     is_complex_dft<descriptor_t>&& std::is_same_v<descriptor_scalar_t<descriptor_t>, data_t>;
diff --git a/include/oneapi/mkl/dft/forward.hpp b/include/oneapi/mkl/dft/forward.hpp
index 353a5a89f..e43c39ce0 100644
--- a/include/oneapi/mkl/dft/forward.hpp
+++ b/include/oneapi/mkl/dft/forward.hpp
@@ -61,9 +61,7 @@ void compute_forward(descriptor_type &desc, sycl::buffer<data_type, 1> &inout_re
 }
 
 //Out-of-place transform
-template <typename descriptor_type, typename input_type, typename output_type,
-          std::enable_if_t<detail::valid_oop_iotypes<descriptor_type, input_type, output_type>,
-                           bool> = true>
+template <typename descriptor_type, typename input_type, typename output_type>
 void compute_forward(descriptor_type &desc, sycl::buffer<input_type, 1> &in,
                      sycl::buffer<output_type, 1> &out) {
     static_assert(detail::valid_compute_arg<descriptor_type, input_type>::value,
@@ -129,9 +127,7 @@ sycl::event compute_forward(descriptor_type &desc, data_type *inout_re, data_typ
 }
 
 //Out-of-place transform
-template <typename descriptor_type, typename input_type, typename output_type,
-          std::enable_if_t<detail::valid_oop_iotypes<descriptor_type, input_type, output_type>,
-                           bool> = true>
+template <typename descriptor_type, typename input_type, typename output_type>
 sycl::event compute_forward(descriptor_type &desc, input_type *in, output_type *out,
                             const std::vector<sycl::event> &dependencies = {}) {
     static_assert(detail::valid_compute_arg<descriptor_type, input_type>::value,