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.skipExternalRttiDefinition = skipExternalRttiDefinition;
112 options.applyTBAA = config.AliasAnalysis;
113 options.forceUnifiedTBAATree = useOldAliasTags;
114 options.typeDescriptorsRenamedForAssembly =
115 !disableCompilerGeneratedNamesConversion;
116 options.ComplexRange = config.ComplexRange;
117 addPassConditionally(pm, disableFirToLlvmIr,
118 [&]() { return fir::createFIRToLLVMPass(options); });
119 // The dialect conversion framework may leave dead unrealized_conversion_cast
120 // ops behind, so run reconcile-unrealized-casts to clean them up.
121 addPassConditionally(pm, disableFirToLlvmIr, [&]() {
122 return mlir::createReconcileUnrealizedCastsPass();
123 });
124}
125
126void addLLVMDialectToLLVMPass(mlir::PassManager &pm,
127 llvm::raw_ostream &output) {
128 addPassConditionally(pm, disableLlvmIrToLlvm, [&]() {
129 return fir::createLLVMDialectToLLVMPass(output);
130 });
131}
132
133void addBoxedProcedurePass(mlir::PassManager &pm) {
134 addPassConditionally(pm, disableBoxedProcedureRewrite,
135 [&]() { return fir::createBoxedProcedurePass(); });
136}
137
138void addExternalNameConversionPass(mlir::PassManager &pm,
139 bool appendUnderscore) {
140 addPassConditionally(pm, disableExternalNameConversion, [&]() {
141 return fir::createExternalNameConversion({appendUnderscore});
142 });
143}
144
145void addCompilerGeneratedNamesConversionPass(mlir::PassManager &pm) {
146 addPassConditionally(pm, disableCompilerGeneratedNamesConversion, [&]() {
147 return fir::createCompilerGeneratedNamesConversion();
148 });
149}
150
151// Use inliner extension point callback to register the default inliner pass.
152void registerDefaultInlinerPass(MLIRToLLVMPassPipelineConfig &config) {
153 config.registerFIRInlinerCallback(
154 [](mlir::PassManager &pm, llvm::OptimizationLevel level) {
155 llvm::StringMap<mlir::OpPassManager> pipelines;
156 // The default inliner pass adds the canonicalizer pass with the default
157 // configuration.
158 pm.addPass(mlir::createInlinerPass(
159 pipelines, addCanonicalizerPassWithoutRegionSimplification));
160 });
161}
162
163/// Create a pass pipeline for running default optimization passes for
164/// incremental conversion of FIR.
165///
166/// \param pm - MLIR pass manager that will hold the pipeline definition
167void createDefaultFIROptimizerPassPipeline(mlir::PassManager &pm,
168 MLIRToLLVMPassPipelineConfig &pc) {
169 // Early Optimizer EP Callback
170 pc.invokeFIROptEarlyEPCallbacks(pm, pc.OptLevel);
171
172 // simplify the IR
173 mlir::GreedyRewriteConfig config;
174 config.setRegionSimplificationLevel(
175 mlir::GreedySimplifyRegionLevel::Disabled);
176 pm.addPass(mlir::createCSEPass());
177 fir::addAVC(pm, pc.OptLevel);
178 addNestedPassToAllTopLevelOperations<PassConstructor>(
179 pm, fir::createCharacterConversion);
180 pm.addPass(mlir::createCanonicalizerPass(config));
181 pm.addPass(fir::createSimplifyRegionLite());
182 if (pc.OptLevel.isOptimizingForSpeed()) {
183 // These passes may increase code size.
184 pm.addPass(fir::createSimplifyIntrinsics());
185 pm.addPass(fir::createAlgebraicSimplificationPass(config));
186 if (enableConstantArgumentGlobalisation)
187 pm.addPass(fir::createConstantArgumentGlobalisationOpt());
188 }
189
190 if (pc.LoopVersioning)
191 pm.addPass(fir::createLoopVersioning());
192
193 pm.addPass(mlir::createCSEPass());
194
195 if (pc.StackArrays)
196 pm.addPass(fir::createStackArrays());
197 else
198 fir::addMemoryAllocationOpt(pm);
199
200 // FIR Inliner Callback
201 pc.invokeFIRInlinerCallback(pm, pc.OptLevel);
202
203 pm.addPass(fir::createSimplifyRegionLite());
204 pm.addPass(mlir::createCSEPass());
205
206 // Polymorphic types
207 pm.addPass(fir::createPolymorphicOpConversion());
208 pm.addPass(fir::createAssumedRankOpConversion());
209
210 // Optimize redundant array repacking operations,
211 // if the source is known to be contiguous.
212 if (pc.OptLevel.isOptimizingForSpeed())
213 pm.addPass(fir::createOptimizeArrayRepacking());
214 pm.addPass(fir::createLowerRepackArraysPass());
215 // Expand FIR operations that may use SCF dialect for their
216 // implementation. This is a mandatory pass.
217 pm.addPass(fir::createSimplifyFIROperations(
218 {/*preferInlineImplementation=*/pc.OptLevel.isOptimizingForSpeed()}));
219
220 if (pc.AliasAnalysis && !disableFirAliasTags && !useOldAliasTags)
221 pm.addPass(fir::createAddAliasTags());
222
223 addNestedPassToAllTopLevelOperations<PassConstructor>(
224 pm, fir::createStackReclaim);
225 // convert control flow to CFG form
226 fir::addCfgConversionPass(pm, pc);
227 pm.addPass(mlir::createSCFToControlFlowPass());
228
229 pm.addPass(mlir::createCanonicalizerPass(config));
230 pm.addPass(fir::createSimplifyRegionLite());
231 pm.addPass(mlir::createCSEPass());
232
233 if (pc.OptLevel.isOptimizingForSpeed())
234 pm.addPass(fir::createSetRuntimeCallAttributes());
235
236 // Last Optimizer EP Callback
237 pc.invokeFIROptLastEPCallbacks(pm, pc.OptLevel);
238}
239
240/// Create a pass pipeline for lowering from HLFIR to FIR
241///
242/// \param pm - MLIR pass manager that will hold the pipeline definition
243/// \param optLevel - optimization level used for creating FIR optimization
244/// passes pipeline
245void createHLFIRToFIRPassPipeline(mlir::PassManager &pm, bool enableOpenMP,
246 llvm::OptimizationLevel optLevel) {
247 if (optLevel.isOptimizingForSpeed()) {
248 addCanonicalizerPassWithoutRegionSimplification(pm);
249 addNestedPassToAllTopLevelOperations<PassConstructor>(
250 pm, hlfir::createSimplifyHLFIRIntrinsics);
251 }
252 addNestedPassToAllTopLevelOperations<PassConstructor>(
253 pm, hlfir::createInlineElementals);
254 if (optLevel.isOptimizingForSpeed()) {
255 addCanonicalizerPassWithoutRegionSimplification(pm);
256 pm.addPass(mlir::createCSEPass());
257 // Run SimplifyHLFIRIntrinsics pass late after CSE,
258 // and allow introducing operations with new side effects.
259 addNestedPassToAllTopLevelOperations<PassConstructor>(pm, []() {
260 return hlfir::createSimplifyHLFIRIntrinsics(
261 {/*allowNewSideEffects=*/true});
262 });
263 addNestedPassToAllTopLevelOperations<PassConstructor>(
264 pm, hlfir::createPropagateFortranVariableAttributes);
265 addNestedPassToAllTopLevelOperations<PassConstructor>(
266 pm, hlfir::createOptimizedBufferization);
267 addNestedPassToAllTopLevelOperations<PassConstructor>(
268 pm, hlfir::createInlineHLFIRAssign);
269
270 if (optLevel == llvm::OptimizationLevel::O3) {
271 addNestedPassToAllTopLevelOperations<PassConstructor>(
272 pm, hlfir::createInlineHLFIRCopyIn);
273 }
274 }
275 pm.addPass(hlfir::createLowerHLFIROrderedAssignments());
276 pm.addPass(hlfir::createLowerHLFIRIntrinsics());
277
278 hlfir::BufferizeHLFIROptions bufferizeOptions;
279 // For opt-for-speed, avoid running any of the loops resulting
280 // from hlfir.elemental lowering, if the result is an empty array.
281 // This helps to avoid long running loops for elementals with
282 // shapes like (0, HUGE).
283 if (optLevel.isOptimizingForSpeed())
284 bufferizeOptions.optimizeEmptyElementals = true;
285 pm.addPass(hlfir::createBufferizeHLFIR(bufferizeOptions));
286 // Run hlfir.assign inlining again after BufferizeHLFIR,
287 // because the latter may introduce new hlfir.assign operations,
288 // e.g. for copying an array into a temporary due to
289 // hlfir.associate.
290 // TODO: we can remove the previous InlineHLFIRAssign, when
291 // FIR AliasAnalysis is good enough to say that a temporary
292 // array does not alias with any user object.
293 if (optLevel.isOptimizingForSpeed())
294 addNestedPassToAllTopLevelOperations<PassConstructor>(
295 pm, hlfir::createInlineHLFIRAssign);
296 pm.addPass(hlfir::createConvertHLFIRtoFIR());
297 if (enableOpenMP)
298 pm.addPass(flangomp::createLowerWorkshare());
299}
300
301/// Create a pass pipeline for handling certain OpenMP transformations needed
302/// prior to FIR lowering.
303///
304/// WARNING: These passes must be run immediately after the lowering to ensure
305/// that the FIR is correct with respect to OpenMP operations/attributes.
306///
307/// \param pm - MLIR pass manager that will hold the pipeline definition.
308/// \param isTargetDevice - Whether code is being generated for a target device
309/// rather than the host device.
310void createOpenMPFIRPassPipeline(mlir::PassManager &pm,
311 OpenMPFIRPassPipelineOpts opts) {
312 using DoConcurrentMappingKind =
313 Fortran::frontend::CodeGenOptions::DoConcurrentMappingKind;
314
315 if (opts.doConcurrentMappingKind != DoConcurrentMappingKind::DCMK_None)
316 pm.addPass(flangomp::createDoConcurrentConversionPass(
317 opts.doConcurrentMappingKind == DoConcurrentMappingKind::DCMK_Device));
318
319 // The MapsForPrivatizedSymbols pass needs to run before
320 // MapInfoFinalizationPass because the former creates new
321 // MapInfoOp instances, typically for descriptors.
322 // MapInfoFinalizationPass adds MapInfoOp instances for the descriptors
323 // underlying data which is necessary to access the data on the offload
324 // target device.
325 pm.addPass(flangomp::createMapsForPrivatizedSymbolsPass());
326 pm.addPass(flangomp::createMapInfoFinalizationPass());
327 pm.addPass(flangomp::createMarkDeclareTargetPass());
328 pm.addPass(flangomp::createGenericLoopConversionPass());
329 if (opts.isTargetDevice)
330 pm.addPass(flangomp::createFunctionFilteringPass());
331}
332
333void createDebugPasses(mlir::PassManager &pm,
334 llvm::codegenoptions::DebugInfoKind debugLevel,
335 llvm::OptimizationLevel OptLevel,
336 llvm::StringRef inputFilename) {
337 if (debugLevel != llvm::codegenoptions::NoDebugInfo)
338 addDebugInfoPass(pm, debugLevel, OptLevel, inputFilename);
339}
340
341void createDefaultFIRCodeGenPassPipeline(mlir::PassManager &pm,
342 MLIRToLLVMPassPipelineConfig config,
343 llvm::StringRef inputFilename) {
344 fir::addBoxedProcedurePass(pm);
345 addNestedPassToAllTopLevelOperations<PassConstructor>(
346 pm, fir::createAbstractResultOpt);
347 addPassToGPUModuleOperations<PassConstructor>(pm,
348 fir::createAbstractResultOpt);
349 fir::addCodeGenRewritePass(
350 pm, (config.DebugInfo != llvm::codegenoptions::NoDebugInfo));
351 fir::addExternalNameConversionPass(pm, config.Underscoring);
352 fir::createDebugPasses(pm, config.DebugInfo, config.OptLevel, inputFilename);
353 fir::addTargetRewritePass(pm);
354 fir::addCompilerGeneratedNamesConversionPass(pm);
355
356 if (config.VScaleMin != 0)
357 pm.addPass(fir::createVScaleAttr({{config.VScaleMin, config.VScaleMax}}));
358
359 // Add function attributes
360 mlir::LLVM::framePointerKind::FramePointerKind framePointerKind;
361
362 if (config.FramePointerKind == llvm::FramePointerKind::NonLeaf)
363 framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::NonLeaf;
364 else if (config.FramePointerKind == llvm::FramePointerKind::All)
365 framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::All;
366 else if (config.FramePointerKind == llvm::FramePointerKind::Reserved)
367 framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::Reserved;
368 else
369 framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::None;
370
371 // TODO: re-enable setNoAlias by default (when optimizing for speed) once
372 // function specialization is fixed.
373 bool setNoAlias = forceNoAlias;
374 bool setNoCapture = config.OptLevel.isOptimizingForSpeed();
375
376 pm.addPass(fir::createFunctionAttr(
377 {framePointerKind, config.InstrumentFunctionEntry,
378 config.InstrumentFunctionExit, config.NoInfsFPMath, config.NoNaNsFPMath,
379 config.ApproxFuncFPMath, config.NoSignedZerosFPMath, config.UnsafeFPMath,
380 config.Reciprocals, config.PreferVectorWidth, /*tuneCPU=*/"",
381 setNoCapture, setNoAlias}));
382
383 if (config.EnableOpenMP) {
384 pm.addNestedPass<mlir::func::FuncOp>(
385 flangomp::createLowerNontemporalPass());
386 }
387
388 fir::addFIRToLLVMPass(pm, config);
389}
390
391/// Create a pass pipeline for lowering from MLIR to LLVM IR
392///
393/// \param pm - MLIR pass manager that will hold the pipeline definition
394/// \param optLevel - optimization level used for creating FIR optimization
395/// passes pipeline
396void createMLIRToLLVMPassPipeline(mlir::PassManager &pm,
397 MLIRToLLVMPassPipelineConfig &config,
398 llvm::StringRef inputFilename) {
399 fir::createHLFIRToFIRPassPipeline(pm, config.EnableOpenMP, config.OptLevel);
400
401 // Add default optimizer pass pipeline.
402 fir::createDefaultFIROptimizerPassPipeline(pm, config);
403
404 // Add codegen pass pipeline.
405 fir::createDefaultFIRCodeGenPassPipeline(pm, config, inputFilename);
406}
407
408} // namespace fir
409

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