438 lines · cpp
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-exception6//7//===----------------------------------------------------------------------===//8 9/// This file defines some utilties to setup FIR pass pipelines. These are10/// 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(true));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::DIEmissionKind84getEmissionKind(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, int32_t dwarfVersion,99 llvm::StringRef splitDwarfFile) {100 fir::AddDebugInfoOptions options;101 options.debugLevel = getEmissionKind(debugLevel);102 options.isOptimized = optLevel != llvm::OptimizationLevel::O0;103 options.inputFilename = inputFilename;104 options.dwarfVersion = dwarfVersion;105 options.splitDwarfFile = splitDwarfFile;106 addPassConditionally(pm, disableDebugInfo,107 [&]() { return fir::createAddDebugInfoPass(options); });108}109 110void addFIRToLLVMPass(mlir::PassManager &pm,111 const MLIRToLLVMPassPipelineConfig &config) {112 fir::FIRToLLVMPassOptions options;113 options.ignoreMissingTypeDescriptors = ignoreMissingTypeDescriptors;114 options.skipExternalRttiDefinition = skipExternalRttiDefinition;115 options.applyTBAA = config.AliasAnalysis;116 options.forceUnifiedTBAATree = useOldAliasTags;117 options.typeDescriptorsRenamedForAssembly =118 !disableCompilerGeneratedNamesConversion;119 options.ComplexRange = config.ComplexRange;120 addPassConditionally(pm, disableFirToLlvmIr,121 [&]() { return fir::createFIRToLLVMPass(options); });122 // The dialect conversion framework may leave dead unrealized_conversion_cast123 // ops behind, so run reconcile-unrealized-casts to clean them up.124 addPassConditionally(pm, disableFirToLlvmIr, [&]() {125 return mlir::createReconcileUnrealizedCastsPass();126 });127}128 129void addLLVMDialectToLLVMPass(mlir::PassManager &pm,130 llvm::raw_ostream &output) {131 addPassConditionally(pm, disableLlvmIrToLlvm, [&]() {132 return fir::createLLVMDialectToLLVMPass(output);133 });134}135 136void addBoxedProcedurePass(mlir::PassManager &pm) {137 addPassConditionally(pm, disableBoxedProcedureRewrite,138 [&]() { return fir::createBoxedProcedurePass(); });139}140 141void addExternalNameConversionPass(mlir::PassManager &pm,142 bool appendUnderscore) {143 addPassConditionally(pm, disableExternalNameConversion, [&]() {144 return fir::createExternalNameConversion({appendUnderscore});145 });146}147 148void addCompilerGeneratedNamesConversionPass(mlir::PassManager &pm) {149 addPassConditionally(pm, disableCompilerGeneratedNamesConversion, [&]() {150 return fir::createCompilerGeneratedNamesConversion();151 });152}153 154// Use inliner extension point callback to register the default inliner pass.155void registerDefaultInlinerPass(MLIRToLLVMPassPipelineConfig &config) {156 config.registerFIRInlinerCallback(157 [](mlir::PassManager &pm, llvm::OptimizationLevel level) {158 llvm::StringMap<mlir::OpPassManager> pipelines;159 // The default inliner pass adds the canonicalizer pass with the default160 // configuration.161 pm.addPass(mlir::createInlinerPass(162 pipelines, addCanonicalizerPassWithoutRegionSimplification));163 });164}165 166/// Create a pass pipeline for running default optimization passes for167/// incremental conversion of FIR.168///169/// \param pm - MLIR pass manager that will hold the pipeline definition170void createDefaultFIROptimizerPassPipeline(mlir::PassManager &pm,171 MLIRToLLVMPassPipelineConfig &pc) {172 // Early Optimizer EP Callback173 pc.invokeFIROptEarlyEPCallbacks(pm, pc.OptLevel);174 175 // simplify the IR176 mlir::GreedyRewriteConfig config;177 config.setRegionSimplificationLevel(178 mlir::GreedySimplifyRegionLevel::Disabled);179 pm.addPass(mlir::createCSEPass());180 fir::addAVC(pm, pc.OptLevel);181 addNestedPassToAllTopLevelOperations<PassConstructor>(182 pm, fir::createCharacterConversion);183 pm.addPass(mlir::createCanonicalizerPass(config));184 pm.addPass(fir::createSimplifyRegionLite());185 if (pc.OptLevel.isOptimizingForSpeed()) {186 // These passes may increase code size.187 pm.addPass(fir::createSimplifyIntrinsics());188 pm.addPass(fir::createAlgebraicSimplificationPass(config));189 if (enableConstantArgumentGlobalisation)190 pm.addPass(fir::createConstantArgumentGlobalisationOpt());191 }192 193 if (pc.LoopVersioning)194 pm.addPass(fir::createLoopVersioning());195 196 pm.addPass(mlir::createCSEPass());197 198 if (pc.StackArrays)199 pm.addPass(fir::createStackArrays());200 else201 fir::addMemoryAllocationOpt(pm);202 203 // FIR Inliner Callback204 pc.invokeFIRInlinerCallback(pm, pc.OptLevel);205 206 pm.addPass(fir::createSimplifyRegionLite());207 pm.addPass(mlir::createCSEPass());208 209 // Polymorphic types210 pm.addPass(fir::createPolymorphicOpConversion());211 pm.addPass(fir::createAssumedRankOpConversion());212 213 // Optimize redundant array repacking operations,214 // if the source is known to be contiguous.215 if (pc.OptLevel.isOptimizingForSpeed())216 pm.addPass(fir::createOptimizeArrayRepacking());217 pm.addPass(fir::createLowerRepackArraysPass());218 // Expand FIR operations that may use SCF dialect for their219 // implementation. This is a mandatory pass.220 pm.addPass(fir::createSimplifyFIROperations(221 {/*preferInlineImplementation=*/pc.OptLevel.isOptimizingForSpeed()}));222 223 addNestedPassToAllTopLevelOperations<PassConstructor>(224 pm, fir::createStackReclaim);225 // convert control flow to CFG form226 fir::addCfgConversionPass(pm, pc);227 pm.addPass(mlir::createSCFToControlFlowPass());228 229 pm.addPass(mlir::createCanonicalizerPass(config));230 pm.addPass(fir::createSimplifyRegionLite());231 if (!pc.SkipConvertComplexPow)232 pm.addPass(fir::createConvertComplexPow());233 pm.addPass(mlir::createCSEPass());234 235 if (pc.OptLevel.isOptimizingForSpeed())236 pm.addPass(fir::createSetRuntimeCallAttributes());237 238 // Last Optimizer EP Callback239 pc.invokeFIROptLastEPCallbacks(pm, pc.OptLevel);240}241 242/// Create a pass pipeline for lowering from HLFIR to FIR243///244/// \param pm - MLIR pass manager that will hold the pipeline definition245/// \param optLevel - optimization level used for creating FIR optimization246/// passes pipeline247void createHLFIRToFIRPassPipeline(mlir::PassManager &pm,248 EnableOpenMP enableOpenMP,249 llvm::OptimizationLevel optLevel) {250 if (optLevel.getSizeLevel() > 0 || optLevel.getSpeedupLevel() > 0) {251 addNestedPassToAllTopLevelOperations<PassConstructor>(252 pm, hlfir::createExpressionSimplification);253 }254 if (optLevel.isOptimizingForSpeed()) {255 addCanonicalizerPassWithoutRegionSimplification(pm);256 addNestedPassToAllTopLevelOperations<PassConstructor>(257 pm, hlfir::createSimplifyHLFIRIntrinsics);258 }259 addNestedPassToAllTopLevelOperations<PassConstructor>(260 pm, hlfir::createInlineElementals);261 if (optLevel.isOptimizingForSpeed()) {262 addCanonicalizerPassWithoutRegionSimplification(pm);263 pm.addPass(mlir::createCSEPass());264 // Run SimplifyHLFIRIntrinsics pass late after CSE,265 // and allow introducing operations with new side effects.266 addNestedPassToAllTopLevelOperations<PassConstructor>(pm, []() {267 return hlfir::createSimplifyHLFIRIntrinsics(268 {/*allowNewSideEffects=*/true});269 });270 addNestedPassToAllTopLevelOperations<PassConstructor>(271 pm, hlfir::createPropagateFortranVariableAttributes);272 addNestedPassToAllTopLevelOperations<PassConstructor>(273 pm, hlfir::createOptimizedBufferization);274 addNestedPassToAllTopLevelOperations<PassConstructor>(275 pm, hlfir::createInlineHLFIRAssign);276 277 if (optLevel == llvm::OptimizationLevel::O3) {278 addNestedPassToAllTopLevelOperations<PassConstructor>(279 pm, hlfir::createInlineHLFIRCopyIn);280 }281 }282 pm.addPass(hlfir::createLowerHLFIROrderedAssignments());283 pm.addPass(hlfir::createLowerHLFIRIntrinsics());284 285 hlfir::BufferizeHLFIROptions bufferizeOptions;286 // For opt-for-speed, avoid running any of the loops resulting287 // from hlfir.elemental lowering, if the result is an empty array.288 // This helps to avoid long running loops for elementals with289 // shapes like (0, HUGE).290 if (optLevel.isOptimizingForSpeed())291 bufferizeOptions.optimizeEmptyElementals = true;292 pm.addPass(hlfir::createBufferizeHLFIR(bufferizeOptions));293 // Run hlfir.assign inlining again after BufferizeHLFIR,294 // because the latter may introduce new hlfir.assign operations,295 // e.g. for copying an array into a temporary due to296 // hlfir.associate.297 // TODO: we can remove the previous InlineHLFIRAssign, when298 // FIR AliasAnalysis is good enough to say that a temporary299 // array does not alias with any user object.300 if (optLevel.isOptimizingForSpeed())301 addNestedPassToAllTopLevelOperations<PassConstructor>(302 pm, hlfir::createInlineHLFIRAssign);303 pm.addPass(hlfir::createConvertHLFIRtoFIR());304 if (enableOpenMP != EnableOpenMP::None) {305 pm.addPass(flangomp::createLowerWorkshare());306 pm.addPass(flangomp::createLowerWorkdistribute());307 }308 if (enableOpenMP == EnableOpenMP::Simd)309 pm.addPass(flangomp::createSimdOnlyPass());310}311 312/// Create a pass pipeline for handling certain OpenMP transformations needed313/// prior to FIR lowering.314///315/// WARNING: These passes must be run immediately after the lowering to ensure316/// that the FIR is correct with respect to OpenMP operations/attributes.317///318/// \param pm - MLIR pass manager that will hold the pipeline definition.319/// \param isTargetDevice - Whether code is being generated for a target device320/// rather than the host device.321void createOpenMPFIRPassPipeline(mlir::PassManager &pm,322 OpenMPFIRPassPipelineOpts opts) {323 using DoConcurrentMappingKind =324 Fortran::frontend::CodeGenOptions::DoConcurrentMappingKind;325 326 if (opts.doConcurrentMappingKind != DoConcurrentMappingKind::DCMK_None)327 pm.addPass(flangomp::createDoConcurrentConversionPass(328 opts.doConcurrentMappingKind == DoConcurrentMappingKind::DCMK_Device));329 330 // The MapsForPrivatizedSymbols and AutomapToTargetDataPass pass need to run331 // before MapInfoFinalizationPass because they create new MapInfoOp332 // instances, typically for descriptors. MapInfoFinalizationPass adds333 // MapInfoOp instances for the descriptors underlying data which is necessary334 // to access the data on the offload target device.335 pm.addPass(flangomp::createMapsForPrivatizedSymbolsPass());336 pm.addPass(flangomp::createAutomapToTargetDataPass());337 pm.addPass(flangomp::createMapInfoFinalizationPass());338 pm.addPass(flangomp::createMarkDeclareTargetPass());339 pm.addPass(flangomp::createGenericLoopConversionPass());340 if (opts.isTargetDevice)341 pm.addPass(flangomp::createFunctionFilteringPass());342}343 344void createDebugPasses(mlir::PassManager &pm,345 llvm::codegenoptions::DebugInfoKind debugLevel,346 llvm::OptimizationLevel OptLevel,347 llvm::StringRef inputFilename, int32_t dwarfVersion,348 llvm::StringRef splitDwarfFile) {349 if (debugLevel != llvm::codegenoptions::NoDebugInfo)350 addDebugInfoPass(pm, debugLevel, OptLevel, inputFilename, dwarfVersion,351 splitDwarfFile);352}353 354void createDefaultFIRCodeGenPassPipeline(mlir::PassManager &pm,355 MLIRToLLVMPassPipelineConfig config,356 llvm::StringRef inputFilename) {357 pm.addPass(fir::createMIFOpConversion());358 fir::addBoxedProcedurePass(pm);359 if (config.OptLevel.isOptimizingForSpeed() && config.AliasAnalysis &&360 !disableFirAliasTags && !useOldAliasTags)361 pm.addPass(fir::createAddAliasTags());362 addNestedPassToAllTopLevelOperations<PassConstructor>(363 pm, fir::createAbstractResultOpt);364 addPassToGPUModuleOperations<PassConstructor>(pm,365 fir::createAbstractResultOpt);366 fir::addCodeGenRewritePass(367 pm, (config.DebugInfo != llvm::codegenoptions::NoDebugInfo));368 fir::addExternalNameConversionPass(pm, config.Underscoring);369 fir::createDebugPasses(pm, config.DebugInfo, config.OptLevel, inputFilename,370 config.DwarfVersion, config.SplitDwarfFile);371 fir::addTargetRewritePass(pm);372 fir::addCompilerGeneratedNamesConversionPass(pm);373 374 if (config.VScaleMin != 0)375 pm.addPass(fir::createVScaleAttr({{config.VScaleMin, config.VScaleMax}}));376 377 // Add function attributes378 mlir::LLVM::framePointerKind::FramePointerKind framePointerKind;379 380 if (config.FramePointerKind == llvm::FramePointerKind::NonLeaf)381 framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::NonLeaf;382 else if (config.FramePointerKind == llvm::FramePointerKind::All)383 framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::All;384 else if (config.FramePointerKind == llvm::FramePointerKind::Reserved)385 framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::Reserved;386 else387 framePointerKind = mlir::LLVM::framePointerKind::FramePointerKind::None;388 389 // TODO: re-enable setNoAlias by default (when optimizing for speed) once390 // function specialization is fixed.391 bool setNoAlias = forceNoAlias;392 bool setNoCapture = config.OptLevel.isOptimizingForSpeed();393 394 pm.addPass(fir::createFunctionAttr(395 {framePointerKind, config.InstrumentFunctionEntry,396 config.InstrumentFunctionExit, config.NoInfsFPMath, config.NoNaNsFPMath,397 config.ApproxFuncFPMath, config.NoSignedZerosFPMath, config.UnsafeFPMath,398 config.Reciprocals, config.PreferVectorWidth, /*tuneCPU=*/"",399 setNoCapture, setNoAlias}));400 401 if (config.EnableOpenMP) {402 pm.addNestedPass<mlir::func::FuncOp>(403 flangomp::createLowerNontemporalPass());404 }405 406 fir::addFIRToLLVMPass(pm, config);407}408 409/// Create a pass pipeline for lowering from MLIR to LLVM IR410///411/// \param pm - MLIR pass manager that will hold the pipeline definition412/// \param optLevel - optimization level used for creating FIR optimization413/// passes pipeline414void createMLIRToLLVMPassPipeline(mlir::PassManager &pm,415 MLIRToLLVMPassPipelineConfig &config,416 llvm::StringRef inputFilename) {417 fir::EnableOpenMP enableOpenMP = fir::EnableOpenMP::None;418 if (config.EnableOpenMP)419 enableOpenMP = fir::EnableOpenMP::Full;420 if (config.EnableOpenMPSimd)421 enableOpenMP = fir::EnableOpenMP::Simd;422 fir::createHLFIRToFIRPassPipeline(pm, enableOpenMP, config.OptLevel);423 424 // Add default optimizer pass pipeline.425 fir::createDefaultFIROptimizerPassPipeline(pm, config);426 427 // Add codegen pass pipeline.428 fir::createDefaultFIRCodeGenPassPipeline(pm, config, inputFilename);429 430 // Run a pass to prepare for translation of delayed privatization in the431 // context of deferred target tasks.432 addPassConditionally(pm, disableFirToLlvmIr, [&]() {433 return mlir::omp::createPrepareForOMPOffloadPrivatizationPass();434 });435}436 437} // namespace fir438