[SYCL][Matrix] Add missing explicit SG size statement (intel#764)

Author: dkhaldi

SYCL/Matrix/element_wise_all_ops_half.cpp

@@ -60,6 +60,9 @@ void matrix_verify_add(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
      auto accA = bufA.get_access<access::mode::read_write>(cgh);
 
 <<<<<<< HEAD
+<<<<<<< HEAD
+=======
+>>>>>>> 7bb961a5c ([SYCL][Matrix] Add missing explicit SG size statement (#764))
      cgh.parallel_for<class add_matrix>(
          r, [accA](nd_item<2> spmd_item) [[intel::reqd_sub_group_size(SG_SZ)]] {
            const auto global_idx = spmd_item.get_global_id(0);
@@ -81,6 +84,7 @@ void matrix_verify_add(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
                                   sg_starty / SG_SZ * TN,
                               N, matrix_layout::row_major);
          }); // parallel for
+<<<<<<< HEAD
 =======
      cgh.parallel_for<class add_matrix>(r, [accA](nd_item<2> spmd_item) {
        const auto global_idx = spmd_item.get_global_id(0);
@@ -103,6 +107,8 @@ void matrix_verify_add(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
                           N, matrix_layout::row_major);
      }); // parallel for
 >>>>>>> 62e420f44 ([SYCL][Matrix] Correct a test case that redefines a class name (#757))
+=======
+>>>>>>> 7bb961a5c ([SYCL][Matrix] Add missing explicit SG size statement (#764))
    }).wait();
   assert_ops_ref<T, M, N>(bufA.get_access<access::mode::read>(), ref);
 }
@@ -116,6 +122,9 @@ void matrix_verify_sub(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
      auto accA = bufA.get_access<access::mode::read_write>(cgh);
 
 <<<<<<< HEAD
+<<<<<<< HEAD
+=======
+>>>>>>> 7bb961a5c ([SYCL][Matrix] Add missing explicit SG size statement (#764))
      cgh.parallel_for<class sub_matrix>(
          r, [accA](nd_item<2> spmd_item) [[intel::reqd_sub_group_size(SG_SZ)]] {
            const auto global_idx = spmd_item.get_global_id(0);
@@ -137,6 +146,7 @@ void matrix_verify_sub(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
                                   sg_starty / SG_SZ * TN,
                               N, matrix_layout::row_major);
          }); // parallel for
+<<<<<<< HEAD
 =======
      cgh.parallel_for<class sub_matrix>(r, [accA](nd_item<2> spmd_item) {
        const auto global_idx = spmd_item.get_global_id(0);
@@ -159,6 +169,8 @@ void matrix_verify_sub(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
                           N, matrix_layout::row_major);
      }); // parallel for
 >>>>>>> 62e420f44 ([SYCL][Matrix] Correct a test case that redefines a class name (#757))
+=======
+>>>>>>> 7bb961a5c ([SYCL][Matrix] Add missing explicit SG size statement (#764))
    }).wait();
   assert_ops_ref<T, M, N>(bufA.get_access<access::mode::read>(), ref);
 }
@@ -172,6 +184,9 @@ void matrix_verify_mul(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
      auto accA = bufA.get_access<access::mode::read_write>(cgh);
 
 <<<<<<< HEAD
+<<<<<<< HEAD
+=======
+>>>>>>> 7bb961a5c ([SYCL][Matrix] Add missing explicit SG size statement (#764))
      cgh.parallel_for<class mul_matrix>(
          r, [accA](nd_item<2> spmd_item) [[intel::reqd_sub_group_size(SG_SZ)]] {
            const auto global_idx = spmd_item.get_global_id(0);
@@ -193,6 +208,7 @@ void matrix_verify_mul(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
                                   sg_starty / SG_SZ * TN,
                               N, matrix_layout::row_major);
          }); // parallel for
+<<<<<<< HEAD
 =======
      cgh.parallel_for<class mul_matrix>(r, [accA](nd_item<2> spmd_item) {
        const auto global_idx = spmd_item.get_global_id(0);
@@ -215,6 +231,8 @@ void matrix_verify_mul(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
                           N, matrix_layout::row_major);
      }); // parallel for
 >>>>>>> 62e420f44 ([SYCL][Matrix] Correct a test case that redefines a class name (#757))
+=======
+>>>>>>> 7bb961a5c ([SYCL][Matrix] Add missing explicit SG size statement (#764))
    }).wait();
   assert_ops_ref<T, M, N>(bufA.get_access<access::mode::read>(), ref);
 }
@@ -228,6 +246,9 @@ void matrix_verify_div(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
      auto accA = bufA.get_access<access::mode::read_write>(cgh);
 
 <<<<<<< HEAD
+<<<<<<< HEAD
+=======
+>>>>>>> 7bb961a5c ([SYCL][Matrix] Add missing explicit SG size statement (#764))
      cgh.parallel_for<class div_matrix>(
          r, [accA](nd_item<2> spmd_item) [[intel::reqd_sub_group_size(SG_SZ)]] {
            const auto global_idx = spmd_item.get_global_id(0);
@@ -249,6 +270,7 @@ void matrix_verify_div(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
                                   sg_starty / SG_SZ * TN,
                               N, matrix_layout::row_major);
          }); // parallel for
+<<<<<<< HEAD
 =======
      cgh.parallel_for<class div_matrix>(r, [accA](nd_item<2> spmd_item) {
        const auto global_idx = spmd_item.get_global_id(0);
@@ -271,6 +293,8 @@ void matrix_verify_div(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
                           N, matrix_layout::row_major);
      }); // parallel for
 >>>>>>> 62e420f44 ([SYCL][Matrix] Correct a test case that redefines a class name (#757))
+=======
+>>>>>>> 7bb961a5c ([SYCL][Matrix] Add missing explicit SG size statement (#764))
    }).wait();
   assert_ops_ref<T, M, N>(bufA.get_access<access::mode::read>(), ref);
 }
@@ -284,6 +308,9 @@ void matrix_verify_logic(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
      auto accA = bufA.get_access<access::mode::read_write>(cgh);
 
 <<<<<<< HEAD
+<<<<<<< HEAD
+=======
+>>>>>>> 7bb961a5c ([SYCL][Matrix] Add missing explicit SG size statement (#764))
      cgh.parallel_for<class logic_matrix>(
          r, [accA](nd_item<2> spmd_item) [[intel::reqd_sub_group_size(SG_SZ)]] {
            const auto global_idx = spmd_item.get_global_id(0);
@@ -314,6 +341,7 @@ void matrix_verify_logic(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
                  }
                  wi_slice_a[i] = val;
                }
+<<<<<<< HEAD
 =======
      cgh.parallel_for<class logic_matrix>(r, [accA](nd_item<2> spmd_item) {
        const auto global_idx = spmd_item.get_global_id(0);
@@ -342,6 +370,8 @@ void matrix_verify_logic(queue q, big_matrix<T, M, N> &A, nd_range<2> &r,
              } else {
                val += 2;
 >>>>>>> 62e420f44 ([SYCL][Matrix] Correct a test case that redefines a class name (#757))
+=======
+>>>>>>> 7bb961a5c ([SYCL][Matrix] Add missing explicit SG size statement (#764))
              }
            }
            joint_matrix_store(sg, sub_a,