[SYCL][Matrix]Align the parameter Group of APIs {joint_matrix_store/joint_matrix_load} with matrix spec (#12041)

[SYCL][Matrix]Align the parameter Group of APIs
{joint_matrix_store/joint_matrix_load} with matrix spec

---------

Signed-off-by: Ni, Wenhui <[email protected]>

Author: wenhuiNi

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

@@ -146,7 +146,7 @@ template <
     std::enable_if_t<std::is_same<S, std::remove_const_t<T>>::value, bool> =
         true>
 inline __SYCL_ALWAYS_INLINE void joint_matrix_load(
-    Group &sg,
+    Group sg,
     joint_matrix<Group, S, use::accumulator, NumRows, NumCols,
                  sycl::ext::oneapi::experimental::matrix::layout::dynamic> &res,
     multi_ptr<T, Space, IsDecorated> src, size_t stride,
@@ -214,7 +214,7 @@ template <
                           std::is_same<std::remove_const_t<T>, float>::value),
                      bool> = true>
 inline __SYCL_ALWAYS_INLINE void
-joint_matrix_load(Group &sg,
+joint_matrix_load(Group sg,
                   joint_matrix<Group, S, Use, NumRows, NumCols, Layout> &res,
                   multi_ptr<T, Space, IsDecorated> src, size_t stride) {
 #if defined(__SYCL_DEVICE_ONLY__)
@@ -253,7 +253,7 @@ joint_matrix_load(Group &sg,
 template <typename Group, typename T, size_t NumRows, size_t NumCols,
           access::address_space Space, access::decorated IsDecorated>
 inline __SYCL_ALWAYS_INLINE void joint_matrix_store(
-    Group &sg,
+    Group sg,
     const joint_matrix<Group, T, use::accumulator, NumRows, NumCols,
                        sycl::ext::oneapi::experimental::matrix::layout::dynamic>
         &src,

sycl/test/matrix/matrix-bfloat16-test.cpp

@@ -59,7 +59,6 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
            const auto sg_startx = global_idx - spmd_item.get_local_id(0);
            const auto sg_starty = global_idy - spmd_item.get_local_id(1);
 
-           sycl::sub_group sg = spmd_item.get_sub_group();
            joint_matrix<sycl::sub_group, bfloat16, use::a, TM, TK,
                         layout::row_major>
                sub_a;
@@ -73,26 +72,27 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
            joint_matrix<sycl::sub_group, float, use::accumulator, TM, TN> sub_c;
 
            joint_matrix_load(
-               sg, sub_c,
+               spmd_item.get_sub_group(), sub_c,
                accC.template get_multi_ptr<sycl::access::decorated::no>() +
                    (sg_startx * TM) * N + sg_starty / SG_SZ * TN,
                N, layout::row_major);
            for (int k = 0; k < K / TK; k += 1) { //
              joint_matrix_load(
-                 sg, sub_a,
+                 spmd_item.get_sub_group(), sub_a,
                  accA.template get_multi_ptr<sycl::access::decorated::no>() +
                      (sg_startx * TM) * K + k * TK,
                  K);
              // Assuming B data is already in VNNI format.
              joint_matrix_load(
-                 sg, sub_b,
+                 spmd_item.get_sub_group(), sub_b,
                  accB.template get_multi_ptr<sycl::access::decorated::no>() +
                      (k * TK / 2) * (N * 2) + sg_starty / SG_SZ * TN * 2,
                  N * 2);
-             joint_matrix_mad(sg, sub_c, sub_a, sub_b, sub_c);
+             joint_matrix_mad(spmd_item.get_sub_group(), sub_c, sub_a, sub_b,
+                              sub_c);
            }
            joint_matrix_store(
-               sg, sub_c,
+               spmd_item.get_sub_group(), sub_c,
                accC.template get_multi_ptr<sycl::access::decorated::no>() +
                    (sg_startx * TM) * N + sg_starty / SG_SZ * TN,
                N, layout::row_major);