fix group_qkv

Author: sunjiweiswift

include/group/epilogue/impl/default_xe.hpp

@@ -100,8 +100,6 @@ class epilogue_t<
         subgroup::msg_type_v<mat_tile_desc, mem_desc_c_t>;
     using matC_payload_t = subgroup::
         mem_payload_t<mem_desc_c_t, mat_tile_desc, msg_type_c, arch_tag>;
-    using matC_payload_t = subgroup::
-        mem_payload_t<mem_desc_c_t, mat_tile_desc, msg_type_c, arch_tag>;
     update_sg_tile_tdesc(g, mem_desc_c);
     matC_t matC;
     matC_payload_t matC_payload(mem_desc_c);

include/subgroup/tile/impl/payload_xe.hpp

@@ -443,7 +443,10 @@ struct mem_payload_t<
     pitch_in_bytes = mem_tdesc.shape.stride * sizeof(dtype);
     width_in_elems = mem_tdesc.shape.x;
     height_in_elems = mem_tdesc.shape.y;
-    payload_bytes = width_in_elems * height_in_elems * sizeof(dtype);
+    payload_bytes = mem_transpose ? (mem_tdesc.shape.x - 1) * pitch_in_bytes +
+            mem_tdesc.shape.y * sizeof(dtype)
+                                  : (mem_tdesc.shape.y - 1) * pitch_in_bytes +
+            mem_tdesc.shape.x * sizeof(dtype);
     uint32_t offset_x = mem_tdesc.coord.x;
     uint32_t offset_y = mem_tdesc.coord.y;
     base_offset = mem_transpose
@@ -464,7 +467,10 @@ struct mem_payload_t<
     uint32_t offset_y = surface_offset_y;
     width_in_elems = surface_width;
     height_in_elems = surface_height;
-    payload_bytes = width_in_elems * height_in_elems * sizeof(dtype);
+    payload_bytes = mem_transpose ? (surface_offset_x - 1) * pitch_in_bytes +
+            surface_offset_y * sizeof(dtype)
+                                  : (surface_offset_y - 1) * pitch_in_bytes +
+            surface_offset_x * sizeof(dtype);
     base_offset = mem_transpose
         ? offset_x * pitch_in_bytes + offset_y * sizeof(dtype)
         : offset_y * pitch_in_bytes + offset_x * sizeof(dtype);
@@ -477,7 +483,11 @@ struct mem_payload_t<
     uint32_t offset_y = mem_tdesc.coord.y;
     width_in_elems = mem_tdesc.shape.x;
     height_in_elems = mem_tdesc.shape.y;
-    payload_bytes = width_in_elems * height_in_elems * sizeof(dtype);
+    payload_bytes = mem_transpose ? (mem_tdesc.shape.x - 1) * pitch_in_bytes +
+            mem_tdesc.shape.y * sizeof(dtype)
+                                  : (mem_tdesc.shape.y - 1) * pitch_in_bytes +
+            mem_tdesc.shape.x * sizeof(dtype);
+
     base_offset = mem_transpose
         ? offset_x * pitch_in_bytes + offset_y * sizeof(dtype)
         : offset_y * pitch_in_bytes + offset_x * sizeof(dtype);
@@ -496,7 +506,9 @@ struct mem_payload_t<
     uint32_t offset_y = surface_offset_y;
     width_in_elems = surface_width;
     height_in_elems = surface_height;
-    payload_bytes = width_in_elems * height_in_elems * sizeof(dtype);
+    payload_bytes = mem_transpose
+        ? (surface_width - 1) * pitch_in_bytes + surface_height * sizeof(dtype)
+        : (surface_height - 1) * pitch_in_bytes + surface_width * sizeof(dtype);
     base_offset = mem_transpose
         ? offset_x * pitch_in_bytes + offset_y * sizeof(dtype)
         : offset_y * pitch_in_bytes + offset_x * sizeof(dtype);

include/subgroup/tile/impl/store_xe.hpp

@@ -282,8 +282,8 @@ tile_store(tile_t& tile, payload_t& payload) {
   using dtype = typename payload_t::dtype;
   static constexpr uint32_t store_len = tile_t::tile_elems;
   static constexpr gpu_arch arch_tag = payload_t::arch_tag;
-
-  if (payload.base_offset <= payload.payload_bytes) {
+  if (payload.base_offset + store_len * sizeof(dtype) <=
+      payload.payload_bytes) {
     using load_store_attr = load_store_attr_t<msg_type::block_1d, arch_tag>;
     static constexpr uint32_t max_store_vec_len =
         load_store_attr::max_store_vec_len;

tests/integration/gemv/int4/main.cpp

@@ -47,7 +47,7 @@ class test_col_major_1 {
   static constexpr mem_layout layout_a = mem_layout::row_major;
   static constexpr mem_layout layout_b = mem_layout::col_major;
   static constexpr mma_engine mma_eng = mma_engine::fpu;
-  static constexpr gpu_arch arch = gpu_arch::XeLpg;
+  static constexpr gpu_arch arch = gpu_arch::XeHpc;
   using data_type_a = fp16;
   using data_type_b = int4x8;
   using data_type_c = fp16;
@@ -108,12 +108,12 @@ int gemm_result_validate(
   bool result = buff_cmp::xetla_buff_cmp(data, other, "gemv validation");
 
 #ifdef UT_DEBUG
-  for (uint32_t i = 0; i < m; i++) {
-    for (uint32_t j = 0; j < n; j++) {
-      std::cout << float(sycl::half(C[i * n + j])) << " ";
-    }
-    std::cout << std::endl;
-  }
+  // for (uint32_t i = 0; i < m; i++) {
+  //   for (uint32_t j = 0; j < n; j++) {
+  //     std::cout << float(sycl::half(C[i * n + j])) << " ";
+  //   }
+  //   std::cout << std::endl;
+  // }
 #endif
   std::cout << (!result ? "FAILED\n" : "PASSED\n");
   return result ? 0 : 1;
@@ -185,12 +185,12 @@ std::vector<data_type_acc_in> dequantize_weight(
     }
   }
 #ifdef UT_DEBUG
-  for (uint32_t i = 0; i < matrix_n; i++) {
-    for (uint32_t j = 0; j < matrix_k; j++) {
-      std::cout << float(sycl::half(b_out[i * matrix_k + j])) << " ";
-    }
-    std::cout << std::endl;
-  }
+  // for (uint32_t i = 0; i < matrix_n; i++) {
+  //   for (uint32_t j = 0; j < matrix_k; j++) {
+  //     std::cout << float(sycl::half(b_out[i * matrix_k + j])) << " ";
+  //   }
+  //   std::cout << std::endl;
+  // }
 #endif
   return b_out;
 }