xref: /llvm-project/flang/lib/Optimizer/Passes/Pipelines.cpp (revision 3bb969f3ebb25037e8eb69c30a5a0dfb5d9d0f51) 
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 
14 namespace fir {
15 
16 template <typename F>
17 void addNestedPassToAllTopLevelOperations(mlir::PassManager &pm, F ctor) {
18   addNestedPassToOps<F, mlir::func::FuncOp, mlir::omp::DeclareReductionOp,
19                      mlir::omp::PrivateClauseOp, fir::GlobalOp>(pm, ctor);
20 }
21 
22 template <typename F>
23 void addPassToGPUModuleOperations(mlir::PassManager &pm, F ctor) {
24   mlir::OpPassManager &nestPM = pm.nest<mlir::gpu::GPUModuleOp>();
25   nestPM.addNestedPass<mlir::func::FuncOp>(ctor());
26   nestPM.addNestedPass<mlir::gpu::GPUFuncOp>(ctor());
27 }
28 
29 template <typename F>
30 void addNestedPassToAllTopLevelOperationsConditionally(
31     mlir::PassManager &pm, llvm::cl::opt<bool> &disabled, F ctor) {
32   if (!disabled)
33     addNestedPassToAllTopLevelOperations<F>(pm, ctor);
34 }
35 
36 void addCanonicalizerPassWithoutRegionSimplification(mlir::OpPassManager &pm) {
37   mlir::GreedyRewriteConfig config;
38   config.enableRegionSimplification = mlir::GreedySimplifyRegionLevel::Disabled;
39   pm.addPass(mlir::createCanonicalizerPass(config));
40 }
41 
42 void addCfgConversionPass(mlir::PassManager &pm,
43                           const MLIRToLLVMPassPipelineConfig &config) {
44   fir::CFGConversionOptions options;
45   if (!config.NSWOnLoopVarInc)
46     options.setNSW = false;
47   addNestedPassToAllTopLevelOperationsConditionally(
48       pm, disableCfgConversion, [&]() { return createCFGConversion(options); });
49 }
50 
51 void addAVC(mlir::PassManager &pm, const llvm::OptimizationLevel &optLevel) {
52   ArrayValueCopyOptions options;
53   options.optimizeConflicts = optLevel.isOptimizingForSpeed();
54   addNestedPassConditionally<mlir::func::FuncOp>(
55       pm, disableFirAvc, [&]() { return createArrayValueCopyPass(options); });
56 }
57 
58 void addMemoryAllocationOpt(mlir::PassManager &pm) {
59   addNestedPassConditionally<mlir::func::FuncOp>(pm, disableFirMao, [&]() {
60     return fir::createMemoryAllocationOpt(
61         {dynamicArrayStackToHeapAllocation, arrayStackAllocationThreshold});
62   });
63 }
64 
65 void addCodeGenRewritePass(mlir::PassManager &pm, bool preserveDeclare) {
66   fir::CodeGenRewriteOptions options;
67   options.preserveDeclare = preserveDeclare;
68   addPassConditionally(pm, disableCodeGenRewrite,
69                        [&]() { return fir::createCodeGenRewrite(options); });
70 }
71 
72 void addTargetRewritePass(mlir::PassManager &pm) {
73   addPassConditionally(pm, disableTargetRewrite,
74                        []() { return fir::createTargetRewritePass(); });
75 }
76 
77 mlir::LLVM::DIEmissionKind
78 getEmissionKind(llvm::codegenoptions::DebugInfoKind kind) {
79   switch (kind) {
80   case llvm::codegenoptions::DebugInfoKind::FullDebugInfo:
81     return mlir::LLVM::DIEmissionKind::Full;
82   case llvm::codegenoptions::DebugInfoKind::DebugLineTablesOnly:
83     return mlir::LLVM::DIEmissionKind::LineTablesOnly;
84   default:
85     return mlir::LLVM::DIEmissionKind::None;
86   }
87 }
88 
89 void addDebugInfoPass(mlir::PassManager &pm,
90                       llvm::codegenoptions::DebugInfoKind debugLevel,
91                       llvm::OptimizationLevel optLevel,
92                       llvm::StringRef inputFilename) {
93   fir::AddDebugInfoOptions options;
94   options.debugLevel = getEmissionKind(debugLevel);
95   options.isOptimized = optLevel != llvm::OptimizationLevel::O0;
96   options.inputFilename = inputFilename;
97   addPassConditionally(pm, disableDebugInfo,
98                        [&]() { return fir::createAddDebugInfoPass(options); });
99 }
100 
101 void addFIRToLLVMPass(mlir::PassManager &pm,
102                       const MLIRToLLVMPassPipelineConfig &config) {
103   fir::FIRToLLVMPassOptions options;
104   options.ignoreMissingTypeDescriptors = ignoreMissingTypeDescriptors;
105   options.applyTBAA = config.AliasAnalysis;
106   options.forceUnifiedTBAATree = useOldAliasTags;
107   options.typeDescriptorsRenamedForAssembly =
108       !disableCompilerGeneratedNamesConversion;
109   addPassConditionally(pm, disableFirToLlvmIr,
110                        [&]() { return fir::createFIRToLLVMPass(options); });
111   // The dialect conversion framework may leave dead unrealized_conversion_cast
112   // ops behind, so run reconcile-unrealized-casts to clean them up.
113   addPassConditionally(pm, disableFirToLlvmIr, [&]() {
114     return mlir::createReconcileUnrealizedCastsPass();
115   });
116 }
117 
118 void addLLVMDialectToLLVMPass(mlir::PassManager &pm,
119                               llvm::raw_ostream &output) {
120   addPassConditionally(pm, disableLlvmIrToLlvm, [&]() {
121     return fir::createLLVMDialectToLLVMPass(output);
122   });
123 }
124 
125 void addBoxedProcedurePass(mlir::PassManager &pm) {
126   addPassConditionally(pm, disableBoxedProcedureRewrite,
127                        [&]() { return fir::createBoxedProcedurePass(); });
128 }
129 
130 void addExternalNameConversionPass(mlir::PassManager &pm,
131                                    bool appendUnderscore) {
132   addPassConditionally(pm, disableExternalNameConversion, [&]() {
133     return fir::createExternalNameConversion({appendUnderscore});
134   });
135 }
136 
137 void addCompilerGeneratedNamesConversionPass(mlir::PassManager &pm) {
138   addPassConditionally(pm, disableCompilerGeneratedNamesConversion, [&]() {
139     return fir::createCompilerGeneratedNamesConversion();
140   });
141 }
142 
143 // Use inliner extension point callback to register the default inliner pass.
144 void registerDefaultInlinerPass(MLIRToLLVMPassPipelineConfig &config) {
145   config.registerFIRInlinerCallback(
146       [](mlir::PassManager &pm, llvm::OptimizationLevel level) {
147         llvm::StringMap<mlir::OpPassManager> pipelines;
148         // The default inliner pass adds the canonicalizer pass with the default
149         // configuration.
150         pm.addPass(mlir::createInlinerPass(
151             pipelines, addCanonicalizerPassWithoutRegionSimplification));
152       });
153 }
154 
155 /// Create a pass pipeline for running default optimization passes for
156 /// incremental conversion of FIR.
157 ///
158 /// \param pm - MLIR pass manager that will hold the pipeline definition
159 void createDefaultFIROptimizerPassPipeline(mlir::PassManager &pm,
160                                            MLIRToLLVMPassPipelineConfig &pc) {
161   // Early Optimizer EP Callback
162   pc.invokeFIROptEarlyEPCallbacks(pm, pc.OptLevel);
163 
164   // simplify the IR
165   mlir::GreedyRewriteConfig config;
166   config.enableRegionSimplification = mlir::GreedySimplifyRegionLevel::Disabled;
167   pm.addPass(mlir::createCSEPass());
168   fir::addAVC(pm, pc.OptLevel);
169   addNestedPassToAllTopLevelOperations<PassConstructor>(
170       pm, fir::createCharacterConversion);
171   pm.addPass(mlir::createCanonicalizerPass(config));
172   pm.addPass(fir::createSimplifyRegionLite());
173   if (pc.OptLevel.isOptimizingForSpeed()) {
174     // These passes may increase code size.
175     pm.addPass(fir::createSimplifyIntrinsics());
176     pm.addPass(fir::createAlgebraicSimplificationPass(config));
177     if (enableConstantArgumentGlobalisation)
178       pm.addPass(fir::createConstantArgumentGlobalisationOpt());
179   }
180 
181   if (pc.LoopVersioning)
182     pm.addPass(fir::createLoopVersioning());
183 
184   pm.addPass(mlir::createCSEPass());
185 
186   if (pc.StackArrays)
187     pm.addPass(fir::createStackArrays());
188   else
189     fir::addMemoryAllocationOpt(pm);
190 
191   // FIR Inliner Callback
192   pc.invokeFIRInlinerCallback(pm, pc.OptLevel);
193 
194   pm.addPass(fir::createSimplifyRegionLite());
195   pm.addPass(mlir::createCSEPass());
196 
197   // Polymorphic types
198   pm.addPass(fir::createPolymorphicOpConversion());
199   pm.addPass(fir::createAssumedRankOpConversion());
200 
201   if (pc.AliasAnalysis && !disableFirAliasTags && !useOldAliasTags)
202     pm.addPass(fir::createAddAliasTags());
203 
204   addNestedPassToAllTopLevelOperations<PassConstructor>(
205       pm, fir::createStackReclaim);
206   // convert control flow to CFG form
207   fir::addCfgConversionPass(pm, pc);
208   pm.addPass(mlir::createConvertSCFToCFPass());
209 
210   pm.addPass(mlir::createCanonicalizerPass(config));
211   pm.addPass(fir::createSimplifyRegionLite());
212   pm.addPass(mlir::createCSEPass());
213 
214   // Last Optimizer EP Callback
215   pc.invokeFIROptLastEPCallbacks(pm, pc.OptLevel);
216 }
217 
218 /// Create a pass pipeline for lowering from HLFIR to FIR
219 ///
220 /// \param pm - MLIR pass manager that will hold the pipeline definition
221 /// \param optLevel - optimization level used for creating FIR optimization
222 ///   passes pipeline
223 void createHLFIRToFIRPassPipeline(mlir::PassManager &pm, bool enableOpenMP,
224                                   llvm::OptimizationLevel optLevel) {
225   if (optLevel.isOptimizingForSpeed()) {
226     addCanonicalizerPassWithoutRegionSimplification(pm);
227     addNestedPassToAllTopLevelOperations<PassConstructor>(
228         pm, hlfir::createSimplifyHLFIRIntrinsics);
229   }
230   addNestedPassToAllTopLevelOperations<PassConstructor>(
231       pm, hlfir::createInlineElementals);
232   if (optLevel.isOptimizingForSpeed()) {
233     addCanonicalizerPassWithoutRegionSimplification(pm);
234     pm.addPass(mlir::createCSEPass());
235     // Run SimplifyHLFIRIntrinsics pass late after CSE,
236     // and allow introducing operations with new side effects.
237     addNestedPassToAllTopLevelOperations<PassConstructor>(pm, []() {
238       return hlfir::createSimplifyHLFIRIntrinsics(
239           {/*allowNewSideEffects=*/true});
240     });
241     addNestedPassToAllTopLevelOperations<PassConstructor>(
242         pm, hlfir::createOptimizedBufferization);
243     addNestedPassToAllTopLevelOperations<PassConstructor>(
244         pm, hlfir::createInlineHLFIRAssign);
245   }
246   pm.addPass(hlfir::createLowerHLFIROrderedAssignments());
247   pm.addPass(hlfir::createLowerHLFIRIntrinsics());
248   pm.addPass(hlfir::createBufferizeHLFIR());
249   // Run hlfir.assign inlining again after BufferizeHLFIR,
250   // because the latter may introduce new hlfir.assign operations,
251   // e.g. for copying an array into a temporary due to
252   // hlfir.associate.
253   // TODO: we can remove the previous InlineHLFIRAssign, when
254   // FIR AliasAnalysis is good enough to say that a temporary
255   // array does not alias with any user object.
256   if (optLevel.isOptimizingForSpeed())
257     addNestedPassToAllTopLevelOperations<PassConstructor>(
258         pm, hlfir::createInlineHLFIRAssign);
259   pm.addPass(hlfir::createConvertHLFIRtoFIR());
260   if (enableOpenMP)
261     pm.addPass(flangomp::createLowerWorkshare());
262 }
263 
264 /// Create a pass pipeline for handling certain OpenMP transformations needed
265 /// prior to FIR lowering.
266 ///
267 /// WARNING: These passes must be run immediately after the lowering to ensure
268 /// that the FIR is correct with respect to OpenMP operations/attributes.
269 ///
270 /// \param pm - MLIR pass manager that will hold the pipeline definition.
271 /// \param isTargetDevice - Whether code is being generated for a target device
272 /// rather than the host device.
273 void createOpenMPFIRPassPipeline(mlir::PassManager &pm, bool isTargetDevice) {
274   pm.addPass(flangomp::createMapInfoFinalizationPass());
275   pm.addPass(flangomp::createMapsForPrivatizedSymbolsPass());
276   pm.addPass(flangomp::createMarkDeclareTargetPass());
277   pm.addPass(flangomp::createGenericLoopConversionPass());
278   if (isTargetDevice)
279     pm.addPass(flangomp::createFunctionFilteringPass());
280 }
281 
282 void createDebugPasses(mlir::PassManager &pm,
283                        llvm::codegenoptions::DebugInfoKind debugLevel,
284                        llvm::OptimizationLevel OptLevel,
285                        llvm::StringRef inputFilename) {
286   if (debugLevel != llvm::codegenoptions::NoDebugInfo)
287     addDebugInfoPass(pm, debugLevel, OptLevel, inputFilename);
288 }
289 
290 void createDefaultFIRCodeGenPassPipeline(mlir::PassManager &pm,
291                                          MLIRToLLVMPassPipelineConfig config,
292                                          llvm::StringRef inputFilename) {
293   fir::addBoxedProcedurePass(pm);
294   addNestedPassToAllTopLevelOperations<PassConstructor>(
295       pm, fir::createAbstractResultOpt);
296   addPassToGPUModuleOperations<PassConstructor>(pm,
297                                                 fir::createAbstractResultOpt);
298   fir::addCodeGenRewritePass(
299       pm, (config.DebugInfo != llvm::codegenoptions::NoDebugInfo));
300   fir::addExternalNameConversionPass(pm, config.Underscoring);
301   fir::createDebugPasses(pm, config.DebugInfo, config.OptLevel, inputFilename);
302   fir::addTargetRewritePass(pm);
303   fir::addCompilerGeneratedNamesConversionPass(pm);
304 
305   if (config.VScaleMin != 0)
306     pm.addPass(fir::createVScaleAttr({{config.VScaleMin, config.VScaleMax}}));
307 
308   // Add function attributes
309   mlir::LLVM::framePointerKind::FramePointerKind framePointerKind;
310 
311   if (config.FramePointerKind == llvm::FramePointerKind::NonLeaf)
312     framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::NonLeaf;
313   else if (config.FramePointerKind == llvm::FramePointerKind::All)
314     framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::All;
315   else
316     framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::None;
317 
318   pm.addPass(fir::createFunctionAttr(
319       {framePointerKind, config.NoInfsFPMath, config.NoNaNsFPMath,
320        config.ApproxFuncFPMath, config.NoSignedZerosFPMath,
321        config.UnsafeFPMath}));
322 
323   fir::addFIRToLLVMPass(pm, config);
324 }
325 
326 /// Create a pass pipeline for lowering from MLIR to LLVM IR
327 ///
328 /// \param pm - MLIR pass manager that will hold the pipeline definition
329 /// \param optLevel - optimization level used for creating FIR optimization
330 ///   passes pipeline
331 void createMLIRToLLVMPassPipeline(mlir::PassManager &pm,
332                                   MLIRToLLVMPassPipelineConfig &config,
333                                   llvm::StringRef inputFilename) {
334   fir::createHLFIRToFIRPassPipeline(pm, config.EnableOpenMP, config.OptLevel);
335 
336   // Add default optimizer pass pipeline.
337   fir::createDefaultFIROptimizerPassPipeline(pm, config);
338 
339   // Add codegen pass pipeline.
340   fir::createDefaultFIRCodeGenPassPipeline(pm, config, inputFilename);
341 }
342 
343 } // namespace fir
344