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 <[email protected]>

Author: s-Nick

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
 

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

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