From ed0f9cdcf56618f43311829c7b21c4654bad5926 Mon Sep 17 00:00:00 2001
From: nscipione <nicolo.scipione@codeplay.com>
Date: Fri, 10 Jan 2025 14:57:56 +0000
Subject: [PATCH 1/3] Update header for SYCL complex implementation and fix
 issue with multi_ptr cast

Newest compiler nightly change include path for header, this change
requires to update to properly check and include sycl complex headers.
multi_ptr cast as it is caused ambigous call to static_cast,
this patch work-around the issue overloading load function using
raw_pointers.

Signed-off-by: nscipione <nicolo.scipione@codeplay.com>
---
 onemath/sycl/blas/include/blas_meta.h         |  6 ++--
 .../blas3/gemm_load_store_complex.hpp         | 14 ++++++++++
 .../blas/src/operations/blas3/gemm_local.hpp  | 28 +++++++++++++++----
 3 files changed, 39 insertions(+), 9 deletions(-)

diff --git a/onemath/sycl/blas/include/blas_meta.h b/onemath/sycl/blas/include/blas_meta.h
index 9e813da..03b19c6 100644
--- a/onemath/sycl/blas/include/blas_meta.h
+++ b/onemath/sycl/blas/include/blas_meta.h
@@ -30,10 +30,10 @@
 #ifdef BLAS_ENABLE_COMPLEX
 #define SYCL_EXT_ONEAPI_COMPLEX
 #include <complex>
-#if __has_include(<ext/oneapi/experimental/complex/complex.hpp>)
-#include <ext/oneapi/experimental/complex/complex.hpp>
+#if __has_include(<sycl/ext/oneapi/experimental/complex/complex.hpp>)
+#include <sycl/ext/oneapi/experimental/complex/complex.hpp>
 #else
-#include <ext/oneapi/experimental/sycl_complex.hpp>
+#include <sycl/ext/oneapi/experimental/sycl_complex.hpp>
 #endif
 #endif
 
diff --git a/onemath/sycl/blas/src/operations/blas3/gemm_load_store_complex.hpp b/onemath/sycl/blas/src/operations/blas3/gemm_load_store_complex.hpp
index 61977df..11934b3 100644
--- a/onemath/sycl/blas/src/operations/blas3/gemm_load_store_complex.hpp
+++ b/onemath/sycl/blas/src/operations/blas3/gemm_load_store_complex.hpp
@@ -116,12 +116,26 @@ class vec_complex {
     m_Data = *(Ptr + Offset * NumElements);
   }
 
+  // Load
+  template <address_t Space, decorated_t DecorateAddress>
+  void load(size_t Offset,
+            const DataT* Ptr) {
+    m_Data = *(Ptr + Offset * NumElements);
+  }
+
   // Store
   template <address_t Space, decorated_t DecorateAddress>
   void store(size_t Offset,
              sycl::multi_ptr<DataT, Space, DecorateAddress> Ptr) const {
     *(Ptr + Offset * NumElements) = m_Data;
   }
+
+  // Store
+  template <address_t Space, decorated_t DecorateAddress>
+  void store(size_t Offset,
+             DataT* Ptr) const {
+    *(Ptr + Offset * NumElements) = m_Data;
+  }
 };
 
 /*! @brief Partial specialization of the Packetize class dedicated to
diff --git a/onemath/sycl/blas/src/operations/blas3/gemm_local.hpp b/onemath/sycl/blas/src/operations/blas3/gemm_local.hpp
index c4c2165..aa5f053 100644
--- a/onemath/sycl/blas/src/operations/blas3/gemm_local.hpp
+++ b/onemath/sycl/blas/src/operations/blas3/gemm_local.hpp
@@ -527,12 +527,28 @@ class Gemm<input_t, output_t, DoubleBuffer, NbcA, NbcB, ClSize, TileType,
       element_t *reg, OutputPointerType out_ptr) {
     vector_out_t out_vec{};
 
-    out_vec.template load<address_t::private_space>(
-        0, sycl::multi_ptr<const element_t, address_t::private_space>(reg));
-    out_vec *= alpha_;
-
-    out_vec.template store<address_t::global_space>(
-        0, sycl::multi_ptr<element_t, address_t::global_space>(out_ptr));
+    if constexpr (std::is_same_v<
+                      element_t,
+                      sycl::ext::oneapi::experimental::complex<float>> ||
+                  std::is_same_v<
+                      element_t,
+                      sycl::ext::oneapi::experimental::complex<double>>) {
+      out_vec.template load<address_t::private_space,
+                            sycl::access::decorated::legacy>(0, reg);
+      out_vec *= alpha_;
+
+    out_vec.template store<address_t::global_space, sycl::access::decorated::legacy>(
+        0, out_ptr);
+    } else {
+      out_vec.template load<address_t::private_space,
+                            sycl::access::decorated::legacy>(
+          0, sycl::multi_ptr<const element_t, address_t::private_space>(reg));
+      out_vec *= alpha_;
+
+      out_vec.template store<address_t::global_space,
+                             sycl::access::decorated::legacy>(
+          0, sycl::multi_ptr<element_t, address_t::global_space>(out_ptr));
+    }
   }
   /*!
    * @brief Store the computed gemm result to the C matrix

From f30d129b09f79fcb5a85f3fa9b58e7064ac3f64b Mon Sep 17 00:00:00 2001
From: nscipione <nicolo.scipione@codeplay.com>
Date: Fri, 10 Jan 2025 15:18:51 +0000
Subject: [PATCH 2/3] Formatting

Signed-off-by: nscipione <nicolo.scipione@codeplay.com>
---
 .../blas/src/operations/blas3/gemm_load_store_complex.hpp   | 6 ++----
 onemath/sycl/blas/src/operations/blas3/gemm_local.hpp       | 4 ++--
 2 files changed, 4 insertions(+), 6 deletions(-)

diff --git a/onemath/sycl/blas/src/operations/blas3/gemm_load_store_complex.hpp b/onemath/sycl/blas/src/operations/blas3/gemm_load_store_complex.hpp
index 11934b3..291fc16 100644
--- a/onemath/sycl/blas/src/operations/blas3/gemm_load_store_complex.hpp
+++ b/onemath/sycl/blas/src/operations/blas3/gemm_load_store_complex.hpp
@@ -118,8 +118,7 @@ class vec_complex {
 
   // Load
   template <address_t Space, decorated_t DecorateAddress>
-  void load(size_t Offset,
-            const DataT* Ptr) {
+  void load(size_t Offset, const DataT *Ptr) {
     m_Data = *(Ptr + Offset * NumElements);
   }
 
@@ -132,8 +131,7 @@ class vec_complex {
 
   // Store
   template <address_t Space, decorated_t DecorateAddress>
-  void store(size_t Offset,
-             DataT* Ptr) const {
+  void store(size_t Offset, DataT *Ptr) const {
     *(Ptr + Offset * NumElements) = m_Data;
   }
 };
diff --git a/onemath/sycl/blas/src/operations/blas3/gemm_local.hpp b/onemath/sycl/blas/src/operations/blas3/gemm_local.hpp
index aa5f053..af01e97 100644
--- a/onemath/sycl/blas/src/operations/blas3/gemm_local.hpp
+++ b/onemath/sycl/blas/src/operations/blas3/gemm_local.hpp
@@ -537,8 +537,8 @@ class Gemm<input_t, output_t, DoubleBuffer, NbcA, NbcB, ClSize, TileType,
                             sycl::access::decorated::legacy>(0, reg);
       out_vec *= alpha_;
 
-    out_vec.template store<address_t::global_space, sycl::access::decorated::legacy>(
-        0, out_ptr);
+      out_vec.template store<address_t::global_space,
+                             sycl::access::decorated::legacy>(0, out_ptr);
     } else {
       out_vec.template load<address_t::private_space,
                             sycl::access::decorated::legacy>(

From f367f9261a84d2995e9de96da6c780534bb2024c Mon Sep 17 00:00:00 2001
From: nscipione <nicolo.scipione@codeplay.com>
Date: Mon, 13 Jan 2025 10:12:48 +0000
Subject: [PATCH 3/3] Adding comment on why the conditional statement is
 necessary

Signed-off-by: nscipione <nicolo.scipione@codeplay.com>
---
 onemath/sycl/blas/src/operations/blas3/gemm_local.hpp | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/onemath/sycl/blas/src/operations/blas3/gemm_local.hpp b/onemath/sycl/blas/src/operations/blas3/gemm_local.hpp
index af01e97..356da25 100644
--- a/onemath/sycl/blas/src/operations/blas3/gemm_local.hpp
+++ b/onemath/sycl/blas/src/operations/blas3/gemm_local.hpp
@@ -527,6 +527,9 @@ class Gemm<input_t, output_t, DoubleBuffer, NbcA, NbcB, ClSize, TileType,
       element_t *reg, OutputPointerType out_ptr) {
     vector_out_t out_vec{};
 
+    // This if-statement is necessary starting from late 2024 nightly, because
+    // an update made casting raw pointers of sycl::complex to multi_ptr
+    // ambiguous.
     if constexpr (std::is_same_v<
                       element_t,
                       sycl::ext::oneapi::experimental::complex<float>> ||