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

[Jianhui Li]

https://github.com/Jianhui-Li created https://github.com/llvm/llvm-project/pull/210539

 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.


>From 2a734232f69e3588ec392626d09c8c0910c818ce Mon Sep 17 00:00:00 2001
From: Jianhui Li <jian.hui.li at intel.com>
Date: Sat, 18 Jul 2026 19:56:27 +0000
Subject: [PATCH] [mlir][gpu] Scope VectorToXeGPU to gpu.module and move
 VectorToSCF to host path

Nest convert-vector-to-xegpu under gpu::GPUModuleOp so it no longer rewrites
host-side vector.transfer ops into xegpu.load/store. Move convert-vector-to-scf
out of the pre-GPU common pipeline and into the post-GPU pipeline so remaining
host transfer ops (and any SCF the pass emits) are lowered before the
vector-to-LLVM conversion.

Co-Authored-By: Claude Opus 4.8 <noreply at anthropic.com>
---
 mlir/lib/Dialect/GPU/Pipelines/CMakeLists.txt        | 1 +
 mlir/lib/Dialect/GPU/Pipelines/GPUToXeVMPipeline.cpp | 3 ++-
 2 files changed, 3 insertions(+), 1 deletion(-)

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());
   {