[XeGPUToXeVM] Add SIMT 4k GEMM e2e test. (#1077)

Author: charithaintc

build_tools/patches/0011-add-mem-copy-support-in-sycl-runtime.patch

@@ -0,0 +1,38 @@
+From 2e17b18af9fe6c10ea4cc87c864ebaaa94d41c18 Mon Sep 17 00:00:00 2001
+From: Charitha Saumya <[email protected]>
+Date: Tue, 20 May 2025 00:12:00 +0000
+Subject: [PATCH] add mem copy support in sycl runtime
+
+---
+ mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp | 9 ++++++++-
+ 1 file changed, 8 insertions(+), 1 deletion(-)
+
+diff --git a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
+index c250340c38fc..5ed9e3871f60 100644
+--- a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
++++ b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
+@@ -149,7 +149,6 @@ static void launchKernel(sycl::queue *queue, sycl::kernel *kernel, size_t gridX,
+ // Wrappers
+
+ extern "C" SYCL_RUNTIME_EXPORT sycl::queue *mgpuStreamCreate() {
+-
+   return catchAll([&]() {
+     sycl::queue *queue =
+         new sycl::queue(getDefaultContext(), getDefaultDevice());
+@@ -168,6 +167,14 @@ mgpuMemAlloc(uint64_t size, sycl::queue *queue, bool isShared) {
+   });
+ }
+
++extern "C" SYCL_RUNTIME_EXPORT void
++mgpuMemcpy(void *dst, void *src, size_t sizeBytes, sycl::queue *queue) {
++  catchAll([&]() {
++    // TODO: Add support for async copy
++    queue->memcpy(dst, src, sizeBytes).wait();
++  });
++}
++
+ extern "C" SYCL_RUNTIME_EXPORT void mgpuMemFree(void *ptr, sycl::queue *queue) {
+   catchAll([&]() {
+     if (ptr) {
+--
+2.34.1

lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp

@@ -274,11 +274,12 @@ class LoadStorePrefetchNdToXeVMPattern : public OpConversionPattern<OpType> {
 
     auto tileW = tdescTy.getDimSize(1);
     auto tileH = tdescTy.getDimSize(0);
-    int32_t vblocks = 1;
-    if (elemBitSize == 16) {
-      vblocks = (tileW + 16 - 1) / 16;
-      tileW = 16;
-    }
+    int32_t vblocks = tdescTy.getArrayLength();
+    // TODO: Why is vblocks calculated like this?
+    // if (elemBitSize == 16) {
+    //   vblocks = (tileW + 16 - 1) / 16;
+    //   tileW = 16;
+    // }
 
     if constexpr (std::is_same_v<OpType, StoreNdOp>) {
       VectorType srcVecTy = cast<VectorType>(op.getValue().getType());

test/CMakeLists.txt

@@ -46,6 +46,12 @@ if(IMEX_ENABLE_VULKAN_RUNNER)
         )
 endif()
 
+if(IMEX_ENABLE_SYCL_RUNTIME)
+    list(APPEND IMEX_TEST_DEPENDS
+        mlir_sycl_runtime
+        )
+endif()
+
 if(IMEX_ENABLE_SYCL_RUNTIME)
     list(APPEND IMEX_TEST_DEPENDS
         sycl-runtime

test/Integration/Dialect/XeGPUToXeVM/gemm_4kx4kx4k_f16_f16_f16_simt.mlir

@@ -7,13 +7,19 @@
         convert-xevm-to-llvm
 	    cse
     )
+    reconcile-unrealized-casts
+    convert-vector-to-scf
     convert-scf-to-cf
+    expand-strided-metadata
+    finalize-memref-to-llvm
     convert-cf-to-llvm
     convert-vector-to-llvm
     convert-arith-to-llvm
-    expand-strided-metadata
-    finalize-memref-to-llvm
+    convert-index-to-llvm
+    convert-func-to-llvm
+    convert-math-to-llvm
     gpu-to-llvm
+    lower-affine
     reconcile-unrealized-casts
     cse
     gpu-module-to-binary)

test/Integration/Dialect/XeGPUToXeVM/xegpu-to-llvm.pp

@@ -40,6 +40,7 @@
 config.substitutions.append(('%imex_runner', config.imex_runner))
 config.substitutions.append(('%python_executable', config.python_executable))
 if config.imex_enable_sycl_runtime:
+    config.substitutions.append(('%mlir_sycl_runtime', config.mlir_sycl_runtime))
     config.substitutions.append(('%sycl_runtime', config.sycl_runtime))
 if config.imex_enable_l0_runtime:
     config.substitutions.append(('%levelzero_runtime', config.levelzero_runtime))

test/lit.cfg.py

@@ -53,6 +53,7 @@ config.mlir_c_runner_utils = os.path.normpath(os.path.join(config.mlir_runner_ut
 if config.enable_vulkan_runner:
     config.vulkan_runtime_wrappers = os.path.normpath(os.path.join(config.mlir_runner_utils_dir, config.shlib_prefix + "vulkan-runtime-wrappers" + config.llvm_shlib_ext))
 if config.imex_enable_sycl_runtime:
+    config.mlir_sycl_runtime = os.path.normpath(os.path.join(config.mlir_runner_utils_dir, config.shlib_prefix + "mlir_sycl_runtime" + config.llvm_shlib_ext))
     config.sycl_runtime = os.path.normpath(os.path.join(config.imex_lib_dir, config.shlib_prefix + "sycl-runtime" + config.llvm_shlib_ext))
 if config.imex_enable_l0_runtime:
     config.levelzero_runtime = os.path.normpath(os.path.join(config.imex_lib_dir, config.shlib_prefix + "level-zero-runtime" + config.llvm_shlib_ext))