1 | //===-- Pipelines.cpp -- FIR pass pipelines ---------------------*- C++ -*-===// |
2 | // |
3 | // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. |
4 | // See https://llvm.org/LICENSE.txt for license information. |
5 | // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception |
6 | // |
7 | //===----------------------------------------------------------------------===// |
8 | |
9 | /// This file defines some utilties to setup FIR pass pipelines. These are |
10 | /// common to flang and the test tools. |
11 | |
12 | #include "flang/Optimizer/Passes/Pipelines.h" |
13 | #include "llvm/Support/CommandLine.h" |
14 | |
15 | /// Force setting the no-alias attribute on fuction arguments when possible. |
16 | static llvm::cl::opt<bool> forceNoAlias("force-no-alias" , llvm::cl::Hidden, |
17 | llvm::cl::init(Val: false)); |
18 | |
19 | namespace fir { |
20 | |
21 | template <typename F> |
22 | void addNestedPassToAllTopLevelOperations(mlir::PassManager &pm, F ctor) { |
23 | addNestedPassToOps<F, mlir::func::FuncOp, mlir::omp::DeclareReductionOp, |
24 | mlir::omp::PrivateClauseOp, fir::GlobalOp>(pm, ctor); |
25 | } |
26 | |
27 | template <typename F> |
28 | void addPassToGPUModuleOperations(mlir::PassManager &pm, F ctor) { |
29 | mlir::OpPassManager &nestPM = pm.nest<mlir::gpu::GPUModuleOp>(); |
30 | nestPM.addNestedPass<mlir::func::FuncOp>(ctor()); |
31 | nestPM.addNestedPass<mlir::gpu::GPUFuncOp>(ctor()); |
32 | } |
33 | |
34 | template <typename F> |
35 | void addNestedPassToAllTopLevelOperationsConditionally( |
36 | mlir::PassManager &pm, llvm::cl::opt<bool> &disabled, F ctor) { |
37 | if (!disabled) |
38 | addNestedPassToAllTopLevelOperations<F>(pm, ctor); |
39 | } |
40 | |
41 | void addCanonicalizerPassWithoutRegionSimplification(mlir::OpPassManager &pm) { |
42 | mlir::GreedyRewriteConfig config; |
43 | config.setRegionSimplificationLevel( |
44 | mlir::GreedySimplifyRegionLevel::Disabled); |
45 | pm.addPass(mlir::createCanonicalizerPass(config)); |
46 | } |
47 | |
48 | void addCfgConversionPass(mlir::PassManager &pm, |
49 | const MLIRToLLVMPassPipelineConfig &config) { |
50 | fir::CFGConversionOptions options; |
51 | if (!config.NSWOnLoopVarInc) |
52 | options.setNSW = false; |
53 | addNestedPassToAllTopLevelOperationsConditionally( |
54 | pm, disableCfgConversion, [&]() { return createCFGConversion(options); }); |
55 | } |
56 | |
57 | void addAVC(mlir::PassManager &pm, const llvm::OptimizationLevel &optLevel) { |
58 | ArrayValueCopyOptions options; |
59 | options.optimizeConflicts = optLevel.isOptimizingForSpeed(); |
60 | addNestedPassConditionally<mlir::func::FuncOp>( |
61 | pm, disableFirAvc, [&]() { return createArrayValueCopyPass(options); }); |
62 | } |
63 | |
64 | void addMemoryAllocationOpt(mlir::PassManager &pm) { |
65 | addNestedPassConditionally<mlir::func::FuncOp>(pm, disableFirMao, [&]() { |
66 | return fir::createMemoryAllocationOpt( |
67 | {dynamicArrayStackToHeapAllocation, arrayStackAllocationThreshold}); |
68 | }); |
69 | } |
70 | |
71 | void addCodeGenRewritePass(mlir::PassManager &pm, bool preserveDeclare) { |
72 | fir::CodeGenRewriteOptions options; |
73 | options.preserveDeclare = preserveDeclare; |
74 | addPassConditionally(pm, disableCodeGenRewrite, |
75 | [&]() { return fir::createCodeGenRewrite(options); }); |
76 | } |
77 | |
78 | void addTargetRewritePass(mlir::PassManager &pm) { |
79 | addPassConditionally(pm, disableTargetRewrite, |
80 | []() { return fir::createTargetRewritePass(); }); |
81 | } |
82 | |
83 | mlir::LLVM::DIEmissionKind |
84 | getEmissionKind(llvm::codegenoptions::DebugInfoKind kind) { |
85 | switch (kind) { |
86 | case llvm::codegenoptions::DebugInfoKind::FullDebugInfo: |
87 | return mlir::LLVM::DIEmissionKind::Full; |
88 | case llvm::codegenoptions::DebugInfoKind::DebugLineTablesOnly: |
89 | return mlir::LLVM::DIEmissionKind::LineTablesOnly; |
90 | default: |
91 | return mlir::LLVM::DIEmissionKind::None; |
92 | } |
93 | } |
94 | |
95 | void addDebugInfoPass(mlir::PassManager &pm, |
96 | llvm::codegenoptions::DebugInfoKind debugLevel, |
97 | llvm::OptimizationLevel optLevel, |
98 | llvm::StringRef inputFilename) { |
99 | fir::AddDebugInfoOptions options; |
100 | options.debugLevel = getEmissionKind(debugLevel); |
101 | options.isOptimized = optLevel != llvm::OptimizationLevel::O0; |
102 | options.inputFilename = inputFilename; |
103 | addPassConditionally(pm, disableDebugInfo, |
104 | [&]() { return fir::createAddDebugInfoPass(options); }); |
105 | } |
106 | |
107 | void addFIRToLLVMPass(mlir::PassManager &pm, |
108 | const MLIRToLLVMPassPipelineConfig &config) { |
109 | fir::FIRToLLVMPassOptions options; |
110 | options.ignoreMissingTypeDescriptors = ignoreMissingTypeDescriptors; |
111 | options.applyTBAA = config.AliasAnalysis; |
112 | options.forceUnifiedTBAATree = useOldAliasTags; |
113 | options.typeDescriptorsRenamedForAssembly = |
114 | !disableCompilerGeneratedNamesConversion; |
115 | addPassConditionally(pm, disableFirToLlvmIr, |
116 | [&]() { return fir::createFIRToLLVMPass(options); }); |
117 | // The dialect conversion framework may leave dead unrealized_conversion_cast |
118 | // ops behind, so run reconcile-unrealized-casts to clean them up. |
119 | addPassConditionally(pm, disableFirToLlvmIr, [&]() { |
120 | return mlir::createReconcileUnrealizedCastsPass(); |
121 | }); |
122 | } |
123 | |
124 | void addLLVMDialectToLLVMPass(mlir::PassManager &pm, |
125 | llvm::raw_ostream &output) { |
126 | addPassConditionally(pm, disableLlvmIrToLlvm, [&]() { |
127 | return fir::createLLVMDialectToLLVMPass(output); |
128 | }); |
129 | } |
130 | |
131 | void addBoxedProcedurePass(mlir::PassManager &pm) { |
132 | addPassConditionally(pm, disableBoxedProcedureRewrite, |
133 | [&]() { return fir::createBoxedProcedurePass(); }); |
134 | } |
135 | |
136 | void addExternalNameConversionPass(mlir::PassManager &pm, |
137 | bool appendUnderscore) { |
138 | addPassConditionally(pm, disableExternalNameConversion, [&]() { |
139 | return fir::createExternalNameConversion({appendUnderscore}); |
140 | }); |
141 | } |
142 | |
143 | void addCompilerGeneratedNamesConversionPass(mlir::PassManager &pm) { |
144 | addPassConditionally(pm, disableCompilerGeneratedNamesConversion, [&]() { |
145 | return fir::createCompilerGeneratedNamesConversion(); |
146 | }); |
147 | } |
148 | |
149 | // Use inliner extension point callback to register the default inliner pass. |
150 | void registerDefaultInlinerPass(MLIRToLLVMPassPipelineConfig &config) { |
151 | config.registerFIRInlinerCallback( |
152 | [](mlir::PassManager &pm, llvm::OptimizationLevel level) { |
153 | llvm::StringMap<mlir::OpPassManager> pipelines; |
154 | // The default inliner pass adds the canonicalizer pass with the default |
155 | // configuration. |
156 | pm.addPass(mlir::createInlinerPass( |
157 | pipelines, addCanonicalizerPassWithoutRegionSimplification)); |
158 | }); |
159 | } |
160 | |
161 | /// Create a pass pipeline for running default optimization passes for |
162 | /// incremental conversion of FIR. |
163 | /// |
164 | /// \param pm - MLIR pass manager that will hold the pipeline definition |
165 | void createDefaultFIROptimizerPassPipeline(mlir::PassManager &pm, |
166 | MLIRToLLVMPassPipelineConfig &pc) { |
167 | // Early Optimizer EP Callback |
168 | pc.invokeFIROptEarlyEPCallbacks(pm, pc.OptLevel); |
169 | |
170 | // simplify the IR |
171 | mlir::GreedyRewriteConfig config; |
172 | config.setRegionSimplificationLevel( |
173 | mlir::GreedySimplifyRegionLevel::Disabled); |
174 | pm.addPass(mlir::createCSEPass()); |
175 | fir::addAVC(pm, pc.OptLevel); |
176 | addNestedPassToAllTopLevelOperations<PassConstructor>( |
177 | pm, fir::createCharacterConversion); |
178 | pm.addPass(mlir::createCanonicalizerPass(config)); |
179 | pm.addPass(fir::createSimplifyRegionLite()); |
180 | if (pc.OptLevel.isOptimizingForSpeed()) { |
181 | // These passes may increase code size. |
182 | pm.addPass(fir::createSimplifyIntrinsics()); |
183 | pm.addPass(fir::createAlgebraicSimplificationPass(config)); |
184 | if (enableConstantArgumentGlobalisation) |
185 | pm.addPass(fir::createConstantArgumentGlobalisationOpt()); |
186 | } |
187 | |
188 | if (pc.LoopVersioning) |
189 | pm.addPass(fir::createLoopVersioning()); |
190 | |
191 | pm.addPass(mlir::createCSEPass()); |
192 | |
193 | if (pc.StackArrays) |
194 | pm.addPass(fir::createStackArrays()); |
195 | else |
196 | fir::addMemoryAllocationOpt(pm); |
197 | |
198 | // FIR Inliner Callback |
199 | pc.invokeFIRInlinerCallback(pm, pc.OptLevel); |
200 | |
201 | pm.addPass(fir::createSimplifyRegionLite()); |
202 | pm.addPass(mlir::createCSEPass()); |
203 | |
204 | // Polymorphic types |
205 | pm.addPass(fir::createPolymorphicOpConversion()); |
206 | pm.addPass(fir::createAssumedRankOpConversion()); |
207 | |
208 | pm.addPass(fir::createLowerRepackArraysPass()); |
209 | // Expand FIR operations that may use SCF dialect for their |
210 | // implementation. This is a mandatory pass. |
211 | pm.addPass(fir::createSimplifyFIROperations( |
212 | {/*preferInlineImplementation=*/pc.OptLevel.isOptimizingForSpeed()})); |
213 | |
214 | if (pc.AliasAnalysis && !disableFirAliasTags && !useOldAliasTags) |
215 | pm.addPass(fir::createAddAliasTags()); |
216 | |
217 | addNestedPassToAllTopLevelOperations<PassConstructor>( |
218 | pm, fir::createStackReclaim); |
219 | // convert control flow to CFG form |
220 | fir::addCfgConversionPass(pm, pc); |
221 | pm.addPass(mlir::createSCFToControlFlowPass()); |
222 | |
223 | pm.addPass(mlir::createCanonicalizerPass(config)); |
224 | pm.addPass(fir::createSimplifyRegionLite()); |
225 | pm.addPass(mlir::createCSEPass()); |
226 | |
227 | if (pc.OptLevel.isOptimizingForSpeed()) |
228 | pm.addPass(fir::createSetRuntimeCallAttributes()); |
229 | |
230 | // Last Optimizer EP Callback |
231 | pc.invokeFIROptLastEPCallbacks(pm, pc.OptLevel); |
232 | } |
233 | |
234 | /// Create a pass pipeline for lowering from HLFIR to FIR |
235 | /// |
236 | /// \param pm - MLIR pass manager that will hold the pipeline definition |
237 | /// \param optLevel - optimization level used for creating FIR optimization |
238 | /// passes pipeline |
239 | void createHLFIRToFIRPassPipeline(mlir::PassManager &pm, bool enableOpenMP, |
240 | llvm::OptimizationLevel optLevel) { |
241 | if (optLevel.isOptimizingForSpeed()) { |
242 | addCanonicalizerPassWithoutRegionSimplification(pm); |
243 | addNestedPassToAllTopLevelOperations<PassConstructor>( |
244 | pm, hlfir::createSimplifyHLFIRIntrinsics); |
245 | } |
246 | addNestedPassToAllTopLevelOperations<PassConstructor>( |
247 | pm, hlfir::createInlineElementals); |
248 | if (optLevel.isOptimizingForSpeed()) { |
249 | addCanonicalizerPassWithoutRegionSimplification(pm); |
250 | pm.addPass(mlir::createCSEPass()); |
251 | // Run SimplifyHLFIRIntrinsics pass late after CSE, |
252 | // and allow introducing operations with new side effects. |
253 | addNestedPassToAllTopLevelOperations<PassConstructor>(pm, []() { |
254 | return hlfir::createSimplifyHLFIRIntrinsics( |
255 | {/*allowNewSideEffects=*/true}); |
256 | }); |
257 | addNestedPassToAllTopLevelOperations<PassConstructor>( |
258 | pm, hlfir::createPropagateFortranVariableAttributes); |
259 | addNestedPassToAllTopLevelOperations<PassConstructor>( |
260 | pm, hlfir::createOptimizedBufferization); |
261 | addNestedPassToAllTopLevelOperations<PassConstructor>( |
262 | pm, hlfir::createInlineHLFIRAssign); |
263 | |
264 | if (optLevel == llvm::OptimizationLevel::O3) { |
265 | addNestedPassToAllTopLevelOperations<PassConstructor>( |
266 | pm, hlfir::createInlineHLFIRCopyIn); |
267 | } |
268 | } |
269 | pm.addPass(hlfir::createLowerHLFIROrderedAssignments()); |
270 | pm.addPass(hlfir::createLowerHLFIRIntrinsics()); |
271 | |
272 | hlfir::BufferizeHLFIROptions bufferizeOptions; |
273 | // For opt-for-speed, avoid running any of the loops resulting |
274 | // from hlfir.elemental lowering, if the result is an empty array. |
275 | // This helps to avoid long running loops for elementals with |
276 | // shapes like (0, HUGE). |
277 | if (optLevel.isOptimizingForSpeed()) |
278 | bufferizeOptions.optimizeEmptyElementals = true; |
279 | pm.addPass(hlfir::createBufferizeHLFIR(bufferizeOptions)); |
280 | // Run hlfir.assign inlining again after BufferizeHLFIR, |
281 | // because the latter may introduce new hlfir.assign operations, |
282 | // e.g. for copying an array into a temporary due to |
283 | // hlfir.associate. |
284 | // TODO: we can remove the previous InlineHLFIRAssign, when |
285 | // FIR AliasAnalysis is good enough to say that a temporary |
286 | // array does not alias with any user object. |
287 | if (optLevel.isOptimizingForSpeed()) |
288 | addNestedPassToAllTopLevelOperations<PassConstructor>( |
289 | pm, hlfir::createInlineHLFIRAssign); |
290 | pm.addPass(hlfir::createConvertHLFIRtoFIR()); |
291 | if (enableOpenMP) |
292 | pm.addPass(flangomp::createLowerWorkshare()); |
293 | } |
294 | |
295 | /// Create a pass pipeline for handling certain OpenMP transformations needed |
296 | /// prior to FIR lowering. |
297 | /// |
298 | /// WARNING: These passes must be run immediately after the lowering to ensure |
299 | /// that the FIR is correct with respect to OpenMP operations/attributes. |
300 | /// |
301 | /// \param pm - MLIR pass manager that will hold the pipeline definition. |
302 | /// \param isTargetDevice - Whether code is being generated for a target device |
303 | /// rather than the host device. |
304 | void createOpenMPFIRPassPipeline(mlir::PassManager &pm, |
305 | OpenMPFIRPassPipelineOpts opts) { |
306 | using DoConcurrentMappingKind = |
307 | Fortran::frontend::CodeGenOptions::DoConcurrentMappingKind; |
308 | |
309 | if (opts.doConcurrentMappingKind != DoConcurrentMappingKind::DCMK_None) |
310 | pm.addPass(flangomp::createDoConcurrentConversionPass( |
311 | opts.doConcurrentMappingKind == DoConcurrentMappingKind::DCMK_Device)); |
312 | |
313 | // The MapsForPrivatizedSymbols pass needs to run before |
314 | // MapInfoFinalizationPass because the former creates new |
315 | // MapInfoOp instances, typically for descriptors. |
316 | // MapInfoFinalizationPass adds MapInfoOp instances for the descriptors |
317 | // underlying data which is necessary to access the data on the offload |
318 | // target device. |
319 | pm.addPass(flangomp::createMapsForPrivatizedSymbolsPass()); |
320 | pm.addPass(flangomp::createMapInfoFinalizationPass()); |
321 | pm.addPass(flangomp::createMarkDeclareTargetPass()); |
322 | pm.addPass(flangomp::createGenericLoopConversionPass()); |
323 | if (opts.isTargetDevice) |
324 | pm.addPass(flangomp::createFunctionFilteringPass()); |
325 | } |
326 | |
327 | void createDebugPasses(mlir::PassManager &pm, |
328 | llvm::codegenoptions::DebugInfoKind debugLevel, |
329 | llvm::OptimizationLevel OptLevel, |
330 | llvm::StringRef inputFilename) { |
331 | if (debugLevel != llvm::codegenoptions::NoDebugInfo) |
332 | addDebugInfoPass(pm, debugLevel, OptLevel, inputFilename); |
333 | } |
334 | |
335 | void createDefaultFIRCodeGenPassPipeline(mlir::PassManager &pm, |
336 | MLIRToLLVMPassPipelineConfig config, |
337 | llvm::StringRef inputFilename) { |
338 | fir::addBoxedProcedurePass(pm); |
339 | addNestedPassToAllTopLevelOperations<PassConstructor>( |
340 | pm, fir::createAbstractResultOpt); |
341 | addPassToGPUModuleOperations<PassConstructor>(pm, |
342 | fir::createAbstractResultOpt); |
343 | fir::addCodeGenRewritePass( |
344 | pm, (config.DebugInfo != llvm::codegenoptions::NoDebugInfo)); |
345 | fir::addExternalNameConversionPass(pm, config.Underscoring); |
346 | fir::createDebugPasses(pm, config.DebugInfo, config.OptLevel, inputFilename); |
347 | fir::addTargetRewritePass(pm); |
348 | fir::addCompilerGeneratedNamesConversionPass(pm); |
349 | |
350 | if (config.VScaleMin != 0) |
351 | pm.addPass(fir::createVScaleAttr({{config.VScaleMin, config.VScaleMax}})); |
352 | |
353 | // Add function attributes |
354 | mlir::LLVM::framePointerKind::FramePointerKind framePointerKind; |
355 | |
356 | if (config.FramePointerKind == llvm::FramePointerKind::NonLeaf) |
357 | framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::NonLeaf; |
358 | else if (config.FramePointerKind == llvm::FramePointerKind::All) |
359 | framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::All; |
360 | else |
361 | framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::None; |
362 | |
363 | // TODO: re-enable setNoAlias by default (when optimizing for speed) once |
364 | // function specialization is fixed. |
365 | bool setNoAlias = forceNoAlias; |
366 | bool setNoCapture = config.OptLevel.isOptimizingForSpeed(); |
367 | |
368 | pm.addPass(fir::createFunctionAttr( |
369 | {framePointerKind, config.InstrumentFunctionEntry, |
370 | config.InstrumentFunctionExit, config.NoInfsFPMath, config.NoNaNsFPMath, |
371 | config.ApproxFuncFPMath, config.NoSignedZerosFPMath, config.UnsafeFPMath, |
372 | config.Reciprocals, config.PreferVectorWidth, /*tuneCPU=*/"" , |
373 | setNoCapture, setNoAlias})); |
374 | |
375 | if (config.EnableOpenMP) { |
376 | pm.addNestedPass<mlir::func::FuncOp>( |
377 | flangomp::createLowerNontemporalPass()); |
378 | } |
379 | |
380 | fir::addFIRToLLVMPass(pm, config); |
381 | } |
382 | |
383 | /// Create a pass pipeline for lowering from MLIR to LLVM IR |
384 | /// |
385 | /// \param pm - MLIR pass manager that will hold the pipeline definition |
386 | /// \param optLevel - optimization level used for creating FIR optimization |
387 | /// passes pipeline |
388 | void createMLIRToLLVMPassPipeline(mlir::PassManager &pm, |
389 | MLIRToLLVMPassPipelineConfig &config, |
390 | llvm::StringRef inputFilename) { |
391 | fir::createHLFIRToFIRPassPipeline(pm, config.EnableOpenMP, config.OptLevel); |
392 | |
393 | // Add default optimizer pass pipeline. |
394 | fir::createDefaultFIROptimizerPassPipeline(pm, config); |
395 | |
396 | // Add codegen pass pipeline. |
397 | fir::createDefaultFIRCodeGenPassPipeline(pm, config, inputFilename); |
398 | } |
399 | |
400 | } // namespace fir |
401 | |