Address review comments.

Author: mshahneo

mlir/include/mlir/Dialect/GPU/Pipelines/Passes.h

@@ -1,4 +1,4 @@
-//===- Passes.h - GPU NVVM/XeVM pipeline entry points----------------------===//
+//===- Passes.h - GPU pipeline entry points----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -63,7 +63,6 @@ struct GPUToNVVMPipelineOptions
 // Options for the gpu to xevm pipeline.
 struct GPUToXeVMPipelineOptions
     : public PassPipelineOptions<GPUToXeVMPipelineOptions> {
-  // XeGPU op granularity selection: workgroup | subgroup | lane
   PassOptions::Option<std::string> xegpuOpLevel{
       *this, "xegpu-op-level",
       llvm::cl::desc("Granularity of XeGPU operations to target: workgroup | "
@@ -123,7 +122,7 @@ void buildLowerToNVVMPassPipeline(OpPassManager &pm,
 void buildLowerToXeVMPassPipeline(OpPassManager &pm,
                                   const GPUToXeVMPipelineOptions &options);
 
-/// Register all pipeleines for the `gpu` dialect.
+/// Register all pipelines for the `gpu` dialect.
 void registerGPUToNVVMPipeline();
 void registerGPUToXeVMPipeline();
 

mlir/lib/Dialect/GPU/Pipelines/CMakeLists.txt

@@ -13,7 +13,6 @@ add_mlir_dialect_library(MLIRGPUPipelines
   MLIRLinalgTransforms
   MLIRAffineToStandard
   MLIRGPUToNVVMTransforms
-  MLIRXeGPUToXeVM
   MLIRIndexToLLVM
   MLIRMathToLLVM
   MLIRNVGPUToNVVM

mlir/lib/Dialect/GPU/Pipelines/GPUToXeVMPipeline.cpp

@@ -40,6 +40,7 @@ void buildCommonPassPipeline(
     OpPassManager &pm, const mlir::gpu::GPUToXeVMPipelineOptions &options) {
   // builtin.module scope passes
   pm.addPass(createCSEPass());
+  pm.addPass(createConvertVectorToSCFPass());
   {
     GpuXeVMAttachTargetOptions xevmTargetOptions;
     xevmTargetOptions.moduleMatcher = options.xevmModuleMatcher;
@@ -49,6 +50,8 @@ void buildCommonPassPipeline(
     xevmTargetOptions.cmdOptions = options.cmdOptions;
     pm.addPass(createGpuXeVMAttachTarget(xevmTargetOptions));
   }
+  pm.addNestedPass<gpu::GPUModuleOp>(createLowerAffinePass());
+  pm.addNestedPass<func::FuncOp>(createGpuAsyncRegionPass());
 }
 
 //===----------------------------------------------------------------------===//
@@ -59,7 +62,6 @@ void buildGpuPassPipeline(OpPassManager &pm,
   if (options.xegpuOpLevel == "workgroup") {
     pm.addNestedPass<gpu::GPUModuleOp>(xegpu::createXeGPUWgToSgDistribute());
     pm.addNestedPass<gpu::GPUModuleOp>(createCSEPass());
-    pm.addNestedPass<gpu::GPUModuleOp>(createLowerAffinePass());
     pm.addNestedPass<gpu::GPUModuleOp>(xegpu::createXeGPUBlocking());
     pm.addNestedPass<gpu::GPUModuleOp>(createCanonicalizerPass());
     pm.addNestedPass<gpu::GPUModuleOp>(createCSEPass());
@@ -90,21 +92,17 @@ void buildGpuPassPipeline(OpPassManager &pm,
 //===----------------------------------------------------------------------===//
 void buildHostPostPipeline(OpPassManager &pm,
                            const mlir::gpu::GPUToXeVMPipelineOptions &options) {
-  pm.addNestedPass<func::FuncOp>(LLVM::createLLVMRequestCWrappersPass());
-  pm.addNestedPass<func::FuncOp>(createGpuAsyncRegionPass());
   pm.addPass(createReconcileUnrealizedCastsPass());
-  pm.addPass(createConvertVectorToSCFPass());
   pm.addPass(createSCFToControlFlowPass());
   pm.addPass(memref::createExpandStridedMetadataPass());
-  pm.addPass(createFinalizeMemRefToLLVMConversionPass());
   {
     GpuToLLVMConversionPassOptions gpuToLLVMOptions;
     gpuToLLVMOptions.hostBarePtrCallConv = options.hostBarePtrCallConv;
     gpuToLLVMOptions.kernelBarePtrCallConv = options.kernelBarePtrCallConv;
     pm.addPass(createGpuToLLVMConversionPass(gpuToLLVMOptions));
   }
-  pm.addPass(createConvertToLLVMPass());
   pm.addPass(createLowerAffinePass());
+  pm.addPass(createConvertToLLVMPass());
   pm.addPass(createReconcileUnrealizedCastsPass());
   // gpu-module-to-binary
   {