@@ -35,11 +35,11 @@ using namespace mlir;
3535
3636namespace {
3737// ===----------------------------------------------------------------------===//
38- // Common pipeline
38+ // Pre-GPU common pipeline for both Host and GPU.
3939// ===----------------------------------------------------------------------===//
40- void buildCommonPassPipeline (
40+ void buildPreGPUCommonPassPipeline (
4141 OpPassManager &pm, const mlir::gpu::GPUToXeVMPipelineOptions &options) {
42- // builtin.module scope passes
42+ // builtin.module scope passes.
4343 pm.addPass (createCSEPass ());
4444 pm.addPass (createConvertVectorToSCFPass ());
4545 {
@@ -58,7 +58,7 @@ void buildCommonPassPipeline(
5858// ===----------------------------------------------------------------------===//
5959// GPUModule-specific stuff.
6060// ===----------------------------------------------------------------------===//
61- void buildGpuPassPipeline (OpPassManager &pm,
61+ void buildGPUPassPipeline (OpPassManager &pm,
6262 const mlir::gpu::GPUToXeVMPipelineOptions &options) {
6363 if (options.xegpuOpLevel == " workgroup" ) {
6464 pm.addNestedPass <gpu::GPUModuleOp>(xegpu::createXeGPUWgToSgDistribute ());
@@ -90,10 +90,11 @@ void buildGpuPassPipeline(OpPassManager &pm,
9090}
9191
9292// ===----------------------------------------------------------------------===//
93- // Host Post-GPU pipeline
93+ // Post-GPU pipeline for both Host and GPU.
9494// ===----------------------------------------------------------------------===//
95- void buildHostPostPipeline (OpPassManager &pm,
95+ void buildPostGPUCommonPassPipeline (OpPassManager &pm,
9696 const mlir::gpu::GPUToXeVMPipelineOptions &options) {
97+ // builtin.module scope passes.
9798 pm.addPass (createSCFToControlFlowPass ());
9899 pm.addPass (memref::createExpandStridedMetadataPass ());
99100 {
@@ -117,14 +118,14 @@ void buildHostPostPipeline(OpPassManager &pm,
117118
118119void mlir::gpu::buildLowerToXeVMPassPipeline (
119120 OpPassManager &pm, const GPUToXeVMPipelineOptions &options) {
120- // Common pipelines
121- buildCommonPassPipeline (pm, options);
121+ // Pre-GPU common pipelines.
122+ buildPreGPUCommonPassPipeline (pm, options);
122123
123- // GPUModule-specific stuff
124- buildGpuPassPipeline (pm, options);
124+ // GPUModule-specific stuff.
125+ buildGPUPassPipeline (pm, options);
125126
126- // Host post-GPUModule-specific stuff
127- buildHostPostPipeline (pm, options);
127+ // Post-GPU pipeline for both Host and GPU.
128+ buildPostGPUCommonPassPipeline (pm, options);
128129}
129130
130131void mlir::gpu::registerGPUToXeVMPipeline () {
0 commit comments