align with IPEX master 0921c6332e3d3e357b9849acc0893a63d9b34b4d ca3e7d24329483babdda0ebff3bca0204c15f735 bac2d0d759c483378bbb41138bf1dc3fe6010026 (#241)

Author: airMeng

include/common/core/arch_config.hpp

@@ -28,6 +28,7 @@ namespace gpu::xetla {
 
 template <msg_type message_type, gpu_arch arch_tag>
 struct load_store_attr_t {};
+
 template <>
 struct load_store_attr_t<msg_type::block_2d, gpu_arch::XeHpc> {
   /// HW limitation checks https://gfxspecs.intel.com/Predator/Home/Index/55490
@@ -75,12 +76,19 @@ struct load_store_attr_t<msg_type::block_2d, gpu_arch::XeHpg>
     : public client_load_store_attr_base_t<
           msg_type::block_2d,
           gpu_arch::XeHpg> {};
+
 template <>
 struct load_store_attr_t<msg_type::block_2d, gpu_arch::XeLpg>
     : public client_load_store_attr_base_t<
           msg_type::block_2d,
           gpu_arch::XeLpg> {};
 
+template <gpu_arch arch_tag>
+struct load_store_attr_t<msg_type::block_1d, arch_tag> {
+  static constexpr uint32_t max_load_vec_len = 64;
+  static constexpr uint32_t max_store_vec_len = 64;
+};
+
 template <gpu_arch arch_tag>
 struct mma_attr_t {};
 

include/common/core/base_types.hpp

@@ -232,7 +232,8 @@ concept xetla_matrix_ref = __ESIMD_NS::detail::is_simd_view_type_v<Ta> &&
 
 } // namespace gpu::xetla
 
-#if (__LIBSYCL_MAJOR_VERSION >= 7) && (__LIBSYCL_MINOR_VERSION >= 1)
+#if (__LIBSYCL_MAJOR_VERSION > 7) || \
+    ((__LIBSYCL_MAJOR_VERSION == 7) && (__LIBSYCL_MINOR_VERSION >= 1))
 
 namespace sycl::detail {
 template <typename T>

include/common/core/common.hpp

@@ -20,6 +20,7 @@
 #pragma once
 
 #include <CL/sycl.hpp>
+#include <common/core/common_types.hpp>
 #include <ext/intel/esimd.hpp>
 #include <version.hpp>
 
@@ -70,15 +71,6 @@ __XETLA_API int32_t xetla_get_subdevice_id() {
 }
 
 namespace gpu::xetla {
-
-enum class gpu_arch : uint8_t { XeLpg = 0, XeHpg = 1, XeHpc = 2 };
-inline constexpr bool arch_has_xmx(gpu_arch arch) {
-  return arch >= gpu_arch::XeHpg;
-}
-
-enum class grf_mode : uint8_t { normal = 0, double_grf = 1 };
-
-enum class mem_layout : uint8_t { row_major = 0, col_major = 1 };
 enum class mem_space : uint8_t { global = 0, local = 1 };
 enum class msg_type : uint8_t {
   block_2d = 0,

include/common/core/common_types.hpp

@@ -0,0 +1,35 @@
+/*******************************************************************************
+ * Copyright (c) 2022-2023 Intel Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *******************************************************************************/
+
+/// @file
+/// C++ API
+
+#pragma once
+#include <cstdint>
+
+namespace gpu::xetla {
+enum class gpu_arch : uint8_t { XeLpg = 0, XeHpg = 1, XeHpc = 2 };
+inline constexpr bool arch_has_xmx(gpu_arch arch) {
+  return arch >= gpu_arch::XeHpg;
+}
+inline constexpr bool arch_has_2d_load_store(gpu_arch arch) {
+  return arch >= gpu_arch::XeHpc;
+}
+
+enum class grf_mode : uint8_t { normal = 0, double_grf = 1 };
+
+enum class mem_layout : uint8_t { row_major = 0, col_major = 1 };
+} // namespace gpu::xetla

include/experimental/common/base_types.hpp

@@ -19,8 +19,6 @@
 
 #pragma once
 
-#include <common/common.hpp>
-
 namespace gpu::xetla {
 
 /// @brief xetla 4bits data packed as 8bits data type.

include/experimental/common/common.hpp

@@ -19,4 +19,5 @@
 
 #pragma once
 
-#include <experimental/common/base_types.hpp>
+#include <common/common.hpp>
+#include <experimental/common/base_types.hpp>

include/group/gemm/impl/default_xmx_xe.hpp

@@ -125,6 +125,8 @@ class gemm_t<
 
   /******** set tile  **********/
   static constexpr reg_layout reg_layout_a = reg_layout::tiled;
+
+ public:
   using matA_tile_desc_t = subgroup::tile_desc_t<
       tile_size_x_a,
       tile_size_y_a,
@@ -165,7 +167,6 @@ class gemm_t<
       wg_size_y,
       arch_tag>;
 
- public:
   using matAcc_tile_desc_t = subgroup::tile_desc_t<
       tile_size_x_c,
       tile_size_y_c,

include/group/gemm/impl/unaligned_xmx_xe.hpp

@@ -129,6 +129,8 @@ class gemm_t<
 
   /******** set tile  **********/
   static constexpr reg_layout reg_layout_a = reg_layout::tiled;
+
+ public:
   using matA_tile_desc_t = subgroup::tile_desc_t<
       tile_size_x_a,
       tile_size_y_a,
@@ -214,7 +216,6 @@ class gemm_t<
       wg_size_y,
       arch_tag>;
 
- public:
   using matAcc_tile_desc_t = subgroup::tile_desc_t<
       tile_size_x_c,
       tile_size_y_c,

include/subgroup/tile/impl/load_xe.hpp

@@ -79,7 +79,7 @@ template <
     typename payload_t>
 __XETLA_API typename std::enable_if_t<
     detail::check_load_type<tile_t, payload_t>::is_global_block_2d &&
-    payload_t::arch_tag == gpu_arch::XeHpc>
+    arch_has_2d_load_store(payload_t::arch_tag)>
 tile_load(tile_t& tile, payload_t& payload) {
   using dtype = typename tile_t::dtype;
   using load_dtype = typename payload_t::mem_dtype;
@@ -405,23 +405,37 @@ tile_load(tile_t& tile, payload_t& payload) {
 
   static constexpr uint32_t tile_size_x = tile_t::tile_size_x;
   static constexpr uint32_t scale_factor = payload_t::scale_factor;
-  constexpr uint32_t load_len = tile_size_x / scale_factor;
+  static constexpr uint32_t load_len = tile_size_x / scale_factor;
+  static constexpr gpu_arch arch_tag = payload_t::arch_tag;
+  using load_store_attr = load_store_attr_t<msg_type::block_1d, arch_tag>;
+  static constexpr uint32_t max_load_vec_len =
+      load_store_attr::max_load_vec_len;
 
-  if constexpr (load_len >= 64) {
+  static constexpr uint32_t load_iter_steps = load_len / max_load_vec_len;
+  if constexpr (load_len >= max_load_vec_len) {
 #pragma unroll
-    for (uint32_t i = 0; i < load_len / 64; i++) {
-      uint32_t offset_x = i * 64 * scale_factor;
-      auto reg_sub = tile.reg.xetla_select<64 * scale_factor, 1>(offset_x);
+    for (uint32_t i = 0; i < load_iter_steps; i++) {
+      uint32_t offset_x = i * max_load_vec_len * scale_factor;
+      auto reg_sub =
+          tile.reg.xetla_select<max_load_vec_len * scale_factor, 1>(offset_x);
       uint32_t address_offset = offset_x * sizeof(dtype);
-      reg_sub.xetla_format<load_dtype>() =
-          xetla_load_global<load_dtype, 64, data_size::default_size, L1, L2>(
-              payload.base_ptr, payload.base_offset + address_offset);
+      reg_sub.xetla_format<load_dtype>() = xetla_load_global<
+          load_dtype,
+          max_load_vec_len,
+          data_size::default_size,
+          L1,
+          L2>(payload.base_ptr, payload.base_offset + address_offset);
     }
   }
-  constexpr uint32_t tail_len = load_len % 64;
-  uint32_t tail_offset = load_len / 64 * 64 * scale_factor;
-  detail::process_1d_tail<tail_len, 32, detail::process_flag::load, L1, L2>(
-      tile, payload, tail_offset);
+
+  constexpr uint32_t tail_len = load_len % max_load_vec_len;
+  uint32_t tail_offset = load_iter_steps * max_load_vec_len * scale_factor;
+  detail::process_1d_tail<
+      tail_len,
+      (max_load_vec_len >> 1),
+      detail::process_flag::load,
+      L1,
+      L2>(tile, payload, tail_offset);
 }
 
 /// @brief This function loads data from unaligned-2D memory surface.
@@ -850,21 +864,33 @@ tile_load(tile_t& tile, payload_t& payload) {
   using load_dtype = typename payload_t::mem_dtype;
 
   constexpr uint32_t scale_factor = payload_t::scale_factor;
-  constexpr uint32_t load_len = tile_desc::tile_size_x / scale_factor;
-  if constexpr (load_len >= 64) {
+  static constexpr uint32_t load_len = tile_desc::tile_size_x / scale_factor;
+  static constexpr gpu_arch arch_tag = payload_t::arch_tag;
+  using load_store_attr = load_store_attr_t<msg_type::block_1d, arch_tag>;
+  static constexpr uint32_t max_load_vec_len =
+      load_store_attr::max_load_vec_len;
+
+  static constexpr uint32_t load_iter_steps = load_len / max_load_vec_len;
+
+  if constexpr (load_len >= max_load_vec_len) {
 #pragma unroll
-    for (uint32_t j = 0; j < load_len / 64; j++) {
-      uint32_t offset_x = j * 64 * scale_factor;
-      auto reg_sub = tile.reg.xetla_select<64 * scale_factor, 1>(offset_x);
+    for (uint32_t j = 0; j < load_iter_steps; j++) {
+      uint32_t offset_x = j * max_load_vec_len * scale_factor;
+      auto reg_sub =
+          tile.reg.xetla_select<max_load_vec_len * scale_factor, 1>(offset_x);
       uint32_t address_offset = offset_x * sizeof(dtype);
-      reg_sub.xetla_format<load_dtype>() =
-          xetla_load_local<load_dtype, 64, data_size::default_size>(
-              payload.address + address_offset);
+      reg_sub.xetla_format<load_dtype>() = xetla_load_local<
+          load_dtype,
+          max_load_vec_len,
+          data_size::default_size>(payload.address + address_offset);
     }
   }
-  detail::
-      process_1d_tail<load_len % 64, 32, detail::process_flag::load, L1, L2>(
-          tile, payload, load_len / 64 * 64 * scale_factor);
+  detail::process_1d_tail<
+      load_len % max_load_vec_len,
+      (max_load_vec_len >> 1),
+      detail::process_flag::load,
+      L1,
+      L2>(tile, payload, load_iter_steps * max_load_vec_len * scale_factor);
 }
 
 } // namespace gpu::xetla::subgroup

include/subgroup/tile/impl/mma_xe.hpp

@@ -38,7 +38,7 @@ struct tile_mma_t<
     matA_t_,
     mma_engine::xmx,
     arch_tag_,
-    std::enable_if_t<(arch_tag_ <= gpu_arch::XeHpc)>> {
+    std::enable_if_t<arch_has_xmx(arch_tag_)>> {
   using matA_t = matA_t_;
   using matB_t = matB_t_;
   using matSrc_t = matAcc_src_t_;