From 2c15e2e3eaa7524ef236c74357e60d613fc8ebca Mon Sep 17 00:00:00 2001
From: Dounia Khaldi <dounia.khaldi@intel.com>
Date: Tue, 15 Oct 2024 13:09:54 -0700
Subject: [PATCH 1/2] [SYCL][Joint Matrix] Add more cases in common JM tests
 functions

---
 sycl/test-e2e/Matrix/common.hpp | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/sycl/test-e2e/Matrix/common.hpp b/sycl/test-e2e/Matrix/common.hpp
index 8d2e46011635..90f5508d97cf 100644
--- a/sycl/test-e2e/Matrix/common.hpp
+++ b/sycl/test-e2e/Matrix/common.hpp
@@ -63,8 +63,7 @@ void matrix_multiply_ref(Ta *A, Tb *B, Tc *C, int M, int N, int K,
           if constexpr (std::is_same_v<Ta, bfloat16> &&
                         std::is_same_v<Tc, float>)
             acc += make_fp32(va[i]) * make_fp32(vb[i]);
-          else if constexpr (std::is_same_v<Ta, sycl::half> &&
-                             std::is_same_v<Tc, float>)
+          else if constexpr (std::is_same_v<Ta, sycl::half>)
             acc += (float)va[i] * (float)vb[i];
           else if constexpr (std::is_same_v<Ta, float> &&
                                  std::is_same_v<Tc, float> ||
@@ -135,7 +134,8 @@ void matrix_rand(unsigned int rows, unsigned int cols, T *src, T val) {
 
   for (unsigned int i = 0; i < rows; i++) {
     for (unsigned int j = 0; j < cols; j++) {
-      if constexpr (std::is_same_v<T, bfloat16> || std::is_same_v<T, float> ||
+      if constexpr (std::is_same_v<T, sycl::half> ||
+                    std::is_same_v<T, bfloat16> || std::is_same_v<T, float> ||
                     std::is_same_v<T, double>) {
         src[i * cols + j] = T(fdistr(dev));
       } else if constexpr (std::is_integral_v<T>) {

From 45e7ba6a6a4ec9e2ffabe264db76815cf32218bd Mon Sep 17 00:00:00 2001
From: Dounia Khaldi <dounia.khaldi@intel.com>
Date: Tue, 5 Nov 2024 06:59:43 -0800
Subject: [PATCH 2/2] restructure combinations so it can be reused and
 augmented by similar architectures

---
 sycl/source/detail/device_info.hpp | 13 +++++++------
 1 file changed, 7 insertions(+), 6 deletions(-)

diff --git a/sycl/source/detail/device_info.hpp b/sycl/source/detail/device_info.hpp
index 82a9e0f14270..3232ed39904f 100644
--- a/sycl/source/detail/device_info.hpp
+++ b/sycl/source/detail/device_info.hpp
@@ -838,8 +838,8 @@ struct get_device_info_impl<
       };
     else if ((architecture::intel_gpu_pvc == DeviceArch) ||
              (architecture::intel_gpu_bmg_g21 == DeviceArch) ||
-             (architecture::intel_gpu_lnl_m == DeviceArch))
-      return {
+             (architecture::intel_gpu_lnl_m == DeviceArch)) {
+      std::vector<ext::oneapi::experimental::matrix::combination> pvc_combs = {
           {8, 0, 0, 0, 16, 32, matrix_type::uint8, matrix_type::uint8,
            matrix_type::sint32, matrix_type::sint32},
           {8, 0, 0, 0, 16, 32, matrix_type::uint8, matrix_type::sint8,
@@ -943,10 +943,11 @@ struct get_device_info_impl<
           {8, 0, 0, 0, 16, 8, matrix_type::tf32, matrix_type::tf32,
            matrix_type::fp32, matrix_type::fp32},
       };
-    else if ((architecture::intel_gpu_dg2_g10 == DeviceArch) ||
-             (architecture::intel_gpu_dg2_g11 == DeviceArch) ||
-             (architecture::intel_gpu_dg2_g12 == DeviceArch) ||
-             (architecture::intel_gpu_arl_h == DeviceArch))
+      return pvc_combs;
+    } else if ((architecture::intel_gpu_dg2_g10 == DeviceArch) ||
+               (architecture::intel_gpu_dg2_g11 == DeviceArch) ||
+               (architecture::intel_gpu_dg2_g12 == DeviceArch) ||
+               (architecture::intel_gpu_arl_h == DeviceArch))
       return {
           {8, 0, 0, 0, 8, 32, matrix_type::uint8, matrix_type::uint8,
            matrix_type::sint32, matrix_type::sint32},