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.
16static llvm::cl::opt<bool> forceNoAlias("force-no-alias", llvm::cl::Hidden,
17 llvm::cl::init(Val: false));
18
19namespace fir {
20
21template <typename F>
22void 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
27template <typename F>
28void 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
34template <typename F>
35void addNestedPassToAllTopLevelOperationsConditionally(
36 mlir::PassManager &pm, llvm::cl::opt<bool> &disabled, F ctor) {
37 if (!disabled)
38 addNestedPassToAllTopLevelOperations<F>(pm, ctor);
39}
40
41void addCanonicalizerPassWithoutRegionSimplification(mlir::OpPassManager &pm) {
42 mlir::GreedyRewriteConfig config;
43 config.setRegionSimplificationLevel(
44 mlir::GreedySimplifyRegionLevel::Disabled);
45 pm.addPass(mlir::createCanonicalizerPass(config));
46}
47
48void 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
57void 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
64void addMemoryAllocationOpt(mlir::PassManager &pm) {
65 addNestedPassConditionally<mlir::func::FuncOp>(pm, disableFirMao, [&]() {
66 return fir::createMemoryAllocationOpt(
67 {dynamicArrayStackToHeapAllocation, arrayStackAllocationThreshold});
68 });
69}
70
71void 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
78void addTargetRewritePass(mlir::PassManager &pm) {
79 addPassConditionally(pm, disableTargetRewrite,
80 []() { return fir::createTargetRewritePass(); });
81}
82
83mlir::LLVM::DIEmissionKind
84getEmissionKind(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
95void 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
107void 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
124void addLLVMDialectToLLVMPass(mlir::PassManager &pm,
125 llvm::raw_ostream &output) {
126 addPassConditionally(pm, disableLlvmIrToLlvm, [&]() {
127 return fir::createLLVMDialectToLLVMPass(output);
128 });
129}
130
131void addBoxedProcedurePass(mlir::PassManager &pm) {
132 addPassConditionally(pm, disableBoxedProcedureRewrite,
133 [&]() { return fir::createBoxedProcedurePass(); });
134}
135
136void addExternalNameConversionPass(mlir::PassManager &pm,
137 bool appendUnderscore) {
138 addPassConditionally(pm, disableExternalNameConversion, [&]() {
139 return fir::createExternalNameConversion({appendUnderscore});
140 });
141}
142
143void 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.
150void 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
165void 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
239void 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.
304void 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
327void 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
335void 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
388void 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

source code of flang/lib/Optimizer/Passes/Pipelines.cpp