[Mlir-commits] [mlir] [mlir][gpu][pipeline] Add missing passes in XeGPU to XeVM pipeline (PR #179320)

[Charitha Saumya]

https://github.com/charithaintc created https://github.com/llvm/llvm-project/pull/179320

as in the title. 

>From a27c0a762005a99773920eed313a1402e5c323f8 Mon Sep 17 00:00:00 2001
From: Charitha Saumya <charitha.saumya.gusthinna.waduge at intel.com>
Date: Mon, 2 Feb 2026 20:53:34 +0000
Subject: [PATCH] save

---
 .../GPU/Pipelines/GPUToXeVMPipeline.cpp       | 25 +++++++++++++------
 1 file changed, 18 insertions(+), 7 deletions(-)

diff --git a/mlir/lib/Dialect/GPU/Pipelines/GPUToXeVMPipeline.cpp b/mlir/lib/Dialect/GPU/Pipelines/GPUToXeVMPipeline.cpp
index f7fff8e1fd4cf..3fd20144d3d3f 100644
--- a/mlir/lib/Dialect/GPU/Pipelines/GPUToXeVMPipeline.cpp
+++ b/mlir/lib/Dialect/GPU/Pipelines/GPUToXeVMPipeline.cpp
@@ -26,6 +26,7 @@
 #include "mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h"
 #include "mlir/Dialect/MemRef/Transforms/Passes.h"
 #include "mlir/Dialect/XeGPU/Transforms/Passes.h"
+#include "mlir/IR/BuiltinOps.h"
 #include "mlir/Pass/PassManager.h"
 #include "mlir/Pass/PassOptions.h"
 #include "mlir/Target/LLVM/XeVM/Target.h"
@@ -60,26 +61,36 @@ void buildPreGPUCommonPassPipeline(
 //===----------------------------------------------------------------------===//
 void buildGPUPassPipeline(OpPassManager &pm,
                           const mlir::gpu::GPUToXeVMPipelineOptions &options) {
+  xegpu::XeGPUPropagateLayoutOptions laneLayoutOptions;
+  laneLayoutOptions.layoutKind = "lane";
+  pm.addNestedPass<ModuleOp>(createCSEPass());
+  pm.addNestedPass<ModuleOp>(createGpuXeVMAttachTarget());
   if (options.xegpuOpLevel == "workgroup") {
+    xegpu::XeGPUPropagateLayoutOptions sgLayoutOptions;
+    sgLayoutOptions.layoutKind = "subgroup";
+    pm.addNestedPass<gpu::GPUModuleOp>(
+        xegpu::createXeGPUPropagateLayout(sgLayoutOptions));
     pm.addNestedPass<gpu::GPUModuleOp>(xegpu::createXeGPUWgToSgDistribute());
     pm.addNestedPass<gpu::GPUModuleOp>(createCSEPass());
-    xegpu::XeGPUPropagateLayoutOptions layoutOptions;
-    layoutOptions.layoutKind = "inst";
+    pm.addNestedPass<gpu::GPUModuleOp>(createLowerAffinePass());
+    pm.addNestedPass<gpu::GPUModuleOp>(createCSEPass());
+    xegpu::XeGPUPropagateLayoutOptions instDataOptions;
+    instDataOptions.layoutKind = "inst";
     pm.addNestedPass<gpu::GPUModuleOp>(
-        xegpu::createXeGPUPropagateLayout(layoutOptions));
+        xegpu::createXeGPUPropagateLayout(instDataOptions));
     pm.addNestedPass<gpu::GPUModuleOp>(xegpu::createXeGPUBlocking());
     pm.addNestedPass<gpu::GPUModuleOp>(createCanonicalizerPass());
     pm.addNestedPass<gpu::GPUModuleOp>(createCSEPass());
   }
   if (options.xegpuOpLevel == "subgroup" ||
       options.xegpuOpLevel == "workgroup") {
-    xegpu::XeGPUPropagateLayoutOptions layoutOptions;
-    layoutOptions.layoutKind = "lane";
     pm.addNestedPass<gpu::GPUModuleOp>(
-        xegpu::createXeGPUPropagateLayout(layoutOptions));
+        xegpu::createXeGPUPropagateLayout(laneLayoutOptions));
     pm.addNestedPass<gpu::GPUModuleOp>(xegpu::createXeGPUPeepHoleOptimizer());
+    pm.addNestedPass<gpu::GPUModuleOp>(createCanonicalizerPass());
+    pm.addNestedPass<gpu::GPUModuleOp>(createCSEPass());
     pm.addNestedPass<gpu::GPUModuleOp>(
-        xegpu::createXeGPUPropagateLayout(layoutOptions));
+        xegpu::createXeGPUPropagateLayout(laneLayoutOptions));
     pm.addNestedPass<gpu::GPUModuleOp>(xegpu::createXeGPUSubgroupDistribute());
     pm.addNestedPass<gpu::GPUModuleOp>(createCanonicalizerPass());
     pm.addNestedPass<gpu::GPUModuleOp>(createCSEPass());