resolve review comments

Author: ggojska

sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp

@@ -276,7 +276,6 @@ class wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, Layout,
     };
     floatValue = *ExtractP;
     return __spirv_ConvertBF16ToFINTEL(intStorage);
-
 #else
     throw exception(make_error_code(errc::runtime),
                     "joint matrix is not supported on host.");

sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp

@@ -87,29 +87,30 @@ extern "C" constexpr __spv::MatrixLayout joint_matrix_layout_to_spv(
 
 template <typename Ta, typename Tb, typename Tc, typename Td>
 constexpr uint32_t CalculateMatrixOperand() {
+  unit32_t returnValue = 0x00;
   if constexpr (std::is_same<Ta, sycl::ext::oneapi::bfloat16>::value &&
                 std::is_same<Tb, sycl::ext::oneapi::bfloat16>::value)
-    return static_cast<uint32_t>(
+    returnValue += static_cast<uint32_t>(
         __spv::MatrixOperands::MatrixAAndBBFloat16ComponentsINTEL);
   if constexpr (std::is_same<Tc, sycl::ext::oneapi::bfloat16>::value)
-    return static_cast<uint32_t>(
+    returnValue +=  static_cast<uint32_t>(
         __spv::MatrixOperands::MatrixCBFloat16ComponentsINTEL);
   if constexpr (std::is_same<Td, sycl::ext::oneapi::bfloat16>::value)
-    return static_cast<uint32_t>(
+    returnValue +=  static_cast<uint32_t>(
         __spv::MatrixOperands::MatrixResultBFloat16ComponentsINTEL);
   if constexpr (std::is_signed<Ta>::value && std::is_unsigned<Tb>::value)
-    return static_cast<uint32_t>(
+    returnValue +=  static_cast<uint32_t>(
         __spv::MatrixOperands::MatrixASignedComponentsKHR);
   if constexpr (std::is_unsigned<Ta>::value && std::is_signed<Tb>::value)
-    return static_cast<uint32_t>(
+    returnValue +=  static_cast<uint32_t>(
         __spv::MatrixOperands::MatrixBSignedComponentsKHR);
   if constexpr (std::is_signed<Ta>::value && std::is_signed<Tb>::value) {
-    return static_cast<uint32_t>(
+    returnValue +=  static_cast<uint32_t>(
         __spv::MatrixOperands::MatrixASignedComponentsKHR) +
            static_cast<uint32_t>(
         __spv::MatrixOperands::MatrixBSignedComponentsKHR);
   }
-  return 0;
+  return returnValue;
 }
 
 } // namespace detail

sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16.cpp

@@ -5,9 +5,10 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
+// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2
 // UNSUPPORTED: gpu-intel-dg2
-// REQUIRES: target-spir
 
+// REQUIRES: target-spir
 // REQUIRES: aspect-ext_intel_matrix
 // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943
 

sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_accumulator.cpp

@@ -5,7 +5,10 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
+// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2
+// UNSUPPORTED: gpu-intel-dg2
 
+// REQUIRES: target-spir
 // REQUIRES: aspect-ext_intel_matrix
 // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943
 

sycl/test-e2e/Matrix/SG32/joint_matrix_half.cpp

@@ -1,10 +1,11 @@
-//==-------- joint_matrix_half.cpp  - DPC++ joint_matrix------------ ----==//
+//==------ SG32/joint_matrix_half.cpp  - DPC++ joint_matrix--------- ----==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
+// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2
 // UNSUPPORTED: gpu-intel-dg2
 
 // REQUIRES: target-spir

sycl/test-e2e/Matrix/SG32/joint_matrix_half_accumulator.cpp

@@ -1,10 +1,12 @@
-//==-------- joint_matrix_half.cpp  - DPC++ joint_matrix------------ ----==//
+//==-------SG32/joint_matrix_half_accumulator.cpp  - DPC++ joint_matrix ----==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
+// SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2
+// UNSUPPORTED: gpu-intel-dg2
 
 // REQUIRES: target-spir
 // REQUIRES: aspect-ext_intel_matrix
@@ -15,6 +17,9 @@
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
+// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=2 %{run} %t.out %}
+// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=1 %{run} %t.out %}
+// RUN: %if gpu %{ env IGC_JointMatrixLoadStoreOpt=0 %{run} %t.out %}
 
 #include "common.hpp"