[Mlir-commits] [mlir] [mlir][gpu] Add VectorToXeGPU and adjust VectorToSCF placement in XeVM pipeline (PR #210539)

[llvmlistbot at llvm.org]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-mlir

Author: Jianhui Li (Jianhui-Li)

<details>
<summary>Changes</summary>

 Add convert-vector-to-xegpu (nested under gpu.module) to lower kernel-side vector ops to XeGPU.
    
Move convert-vector-to-scf from the pre-GPU common pipeline into the post-GPU pipeline, so host transfer ops (and any SCF the pass emits) are lowered to loops before the vector-to-LLVM conversion.


---
Full diff: https://github.com/llvm/llvm-project/pull/210539.diff


2 Files Affected:

- (modified) mlir/lib/Dialect/GPU/Pipelines/CMakeLists.txt (+1) 
- (modified) mlir/lib/Dialect/GPU/Pipelines/GPUToXeVMPipeline.cpp (+2-1) 


``````````diff
diff --git a/mlir/lib/Dialect/GPU/Pipelines/CMakeLists.txt b/mlir/lib/Dialect/GPU/Pipelines/CMakeLists.txt
index 293166cbf8b83..3a0cd7ef5c39d 100644
--- a/mlir/lib/Dialect/GPU/Pipelines/CMakeLists.txt
+++ b/mlir/lib/Dialect/GPU/Pipelines/CMakeLists.txt
@@ -30,6 +30,7 @@ add_mlir_dialect_library(MLIRGPUPipelines
   MLIRSCFToControlFlow
   MLIRVectorToLLVMPass
   MLIRVectorToSCF
+  MLIRVectorToXeGPU
   MLIRXeGPUTransforms
   MLIRXeGPUToXeVM
   MLIRXeVMToLLVM
diff --git a/mlir/lib/Dialect/GPU/Pipelines/GPUToXeVMPipeline.cpp b/mlir/lib/Dialect/GPU/Pipelines/GPUToXeVMPipeline.cpp
index 1cbf3bd633633..d0bab5cd74220 100644
--- a/mlir/lib/Dialect/GPU/Pipelines/GPUToXeVMPipeline.cpp
+++ b/mlir/lib/Dialect/GPU/Pipelines/GPUToXeVMPipeline.cpp
@@ -44,7 +44,6 @@ void buildPreGPUCommonPassPipeline(
     OpPassManager &pm, const mlir::gpu::GPUToXeVMPipelineOptions &options) {
   // builtin.module scope passes.
   pm.addPass(createCSEPass());
-  pm.addPass(createConvertVectorToSCFPass());
   {
     GpuXeVMAttachTargetOptions xevmTargetOptions;
     xevmTargetOptions.moduleMatcher = options.xevmModuleMatcher;
@@ -67,6 +66,7 @@ void buildGPUPassPipeline(OpPassManager &pm,
   laneLayoutOptions.indexBitWidth = options.use64bitIndex ? 64 : 32;
   laneLayoutOptions.layoutKind = "lane";
   pm.addNestedPass<ModuleOp>(createCSEPass());
+  pm.addNestedPass<gpu::GPUModuleOp>(createConvertVectorToXeGPU());
   if (options.xegpuOpLevel == "workgroup") {
     xegpu::XeGPUPropagateLayoutOptions sgLayoutOptions;
     sgLayoutOptions.layoutKind = "subgroup";
@@ -154,6 +154,7 @@ void buildGPUPassPipeline(OpPassManager &pm,
 void buildPostGPUCommonPassPipeline(
     OpPassManager &pm, const mlir::gpu::GPUToXeVMPipelineOptions &options) {
   // builtin.module scope passes.
+  pm.addPass(createConvertVectorToSCFPass());
   pm.addPass(createSCFToControlFlowPass());
   pm.addPass(memref::createExpandStridedMetadataPass());
   {

``````````

</details>


https://github.com/llvm/llvm-project/pull/210539