Address review comments.

Author: mshahneo

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

@@ -1,4 +1,4 @@
-//===- Passes.h - GPU 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.

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

@@ -50,7 +50,7 @@ void buildCommonPassPipeline(
     xevmTargetOptions.cmdOptions = options.cmdOptions;
     pm.addPass(createGpuXeVMAttachTarget(xevmTargetOptions));
   }
-  pm.addNestedPass<gpu::GPUModuleOp>(createLowerAffinePass());
+  pm.addPass(createLowerAffinePass());
   pm.addNestedPass<func::FuncOp>(createGpuAsyncRegionPass());
 }
 
@@ -84,14 +84,14 @@ void buildGpuPassPipeline(OpPassManager &pm,
         createConvertGpuOpsToLLVMSPVOps(gpuToLLVMSPVOptions));
   }
   pm.addNestedPass<gpu::GPUModuleOp>(createCSEPass());
+  pm.addPass(createReconcileUnrealizedCastsPass());
 }
 
 //===----------------------------------------------------------------------===//
 // Host Post-GPU pipeline
 //===----------------------------------------------------------------------===//
 void buildHostPostPipeline(OpPassManager &pm,
                            const mlir::gpu::GPUToXeVMPipelineOptions &options) {
-  pm.addPass(createReconcileUnrealizedCastsPass());
   pm.addPass(createSCFToControlFlowPass());
   pm.addPass(memref::createExpandStridedMetadataPass());
   {

mlir/test/Integration/Dialect/XeGPU/WG/simple_gemm.mlir

@@ -26,7 +26,9 @@ module @gemm attributes {gpu.container_module} {
     gpu.memcpy %B_gpu, %B : memref<256x256xf16>, memref<256x256xf16>
     %C_gpu = gpu.alloc () : memref<256x256xf32>
     gpu.memcpy %C_gpu, %C : memref<256x256xf32>, memref<256x256xf32>
-    // NOTE: Here we can't use [8, 64] wi threads following the SG thread layout of [8, 4]. Because runtime will linearize the x dimension first (we need y dimension to be linearized first).
+    // NOTE: Here we can't use [8, 64] wi threads following
+    // the SG thread layout of [8, 4]. Because runtime will linearize
+    // the x dimension first (we need y dimension to be linearized first).
     // So just use linearized thread layout of [512, 1] wi threads.
     gpu.launch_func  @test_kernel::@test_kernel blocks in (%c1, %c1, %c1) threads in (%c512, %c1, %c1) args(%A_gpu : memref<256x256xf16>, %B_gpu : memref<256x256xf16>, %C_gpu : memref<256x256xf32>)
     gpu.wait // Wait for the kernel to finish.