caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[flang][cuda] Change how abstract result pass is scheduled on func.func and gpu.func (#119034)

Browse Source 
Use `pm.nest` to schedule the pass on nested `func.func` and `gpu.func`
in the `gpu.module`.

AbstractResult pass is not meant to run on the whole gpu.module at once.
This commit is contained in:
Valentin Clement (バレンタイン クレメン)
2024-12-09 13:31:27 -08:00
committed by GitHub
parent f15cc6fe2f
commit 1d4b5c161f
6 changed files with 40 additions and 61 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
flang/lib/Optimizer/Passes/Pipelines.cpp
View File

@@ -16,8 +16,13 @@ namespace fir {
void addNestedPassToAllTopLevelOperations(mlir::PassManager &pm,
PassConstructor ctor) {
addNestedPassToOps<mlir::func::FuncOp, mlir::omp::DeclareReductionOp,
mlir::omp::PrivateClauseOp, fir::GlobalOp,
mlir::gpu::GPUModuleOp>(pm, ctor);
mlir::omp::PrivateClauseOp, fir::GlobalOp>(pm, ctor);
}
void addPassToGPUModuleOperations(mlir::PassManager &pm, PassConstructor ctor) {
mlir::OpPassManager &nestPM = pm.nest<mlir::gpu::GPUModuleOp>();
nestPM.addNestedPass<mlir::func::FuncOp>(ctor());
nestPM.addNestedPass<mlir::gpu::GPUFuncOp>(ctor());
}
void addNestedPassToAllTopLevelOperationsConditionally(
@@ -266,6 +271,7 @@ void createDefaultFIRCodeGenPassPipeline(mlir::PassManager &pm,
llvm::StringRef inputFilename) {
fir::addBoxedProcedurePass(pm);
addNestedPassToAllTopLevelOperations(pm, fir::createAbstractResultOpt);
addPassToGPUModuleOperations(pm, fir::createAbstractResultOpt);
fir::addCodeGenRewritePass(
pm, (config.DebugInfo != llvm::codegenoptions::NoDebugInfo));
fir::addExternalNameConversionPass(pm, config.Underscoring);

15
flang/lib/Optimizer/Transforms/AbstractResult.cpp
View File

@@ -460,17 +460,10 @@ public:
const bool shouldBoxResult = this->passResultAsBox.getValue();
mlir::TypeSwitch<mlir::Operation *, void>(op)
.Case<mlir::func::FuncOp, fir::GlobalOp>([&](auto op) {
runOnSpecificOperation(op, shouldBoxResult, patterns, target);
})
.Case<mlir::gpu::GPUModuleOp>([&](auto op) {
auto gpuMod = mlir::dyn_cast<mlir::gpu::GPUModuleOp>(*op);
for (auto funcOp : gpuMod.template getOps<mlir::func::FuncOp>())
runOnSpecificOperation(funcOp, shouldBoxResult, patterns, target);
for (auto gpuFuncOp : gpuMod.template getOps<mlir::gpu::GPUFuncOp>())
runOnSpecificOperation(gpuFuncOp, shouldBoxResult, patterns,
target);
});
.Case<mlir::func::FuncOp, fir::GlobalOp, mlir::gpu::GPUFuncOp>(
[&](auto op) {
runOnSpecificOperation(op, shouldBoxResult, patterns, target);
});
// Convert the calls and, if needed, the ReturnOp in the function body.
target.addLegalDialect<fir::FIROpsDialect, mlir::arith::ArithDialect,

9
flang/test/Driver/bbc-mlir-pass-pipeline.f90
View File

@@ -17,14 +17,12 @@ end program
! CHECK-NEXT: (S) 0 num-cse'd - Number of operations CSE'd
! CHECK-NEXT: (S) 0 num-dce'd - Number of operations DCE'd
! CHECK-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
! CHECK-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
! CHECK-NEXT: 'fir.global' Pipeline
! CHECK-NEXT: CharacterConversion
! CHECK-NEXT: 'func.func' Pipeline
! CHECK-NEXT: ArrayValueCopy
! CHECK-NEXT: CharacterConversion
! CHECK-NEXT: 'gpu.module' Pipeline
! CHECK-NEXT: CharacterConversion
! CHECK-NEXT: 'omp.declare_reduction' Pipeline
! CHECK-NEXT: CharacterConversion
! CHECK-NEXT: 'omp.private' Pipeline
@@ -50,16 +48,13 @@ end program
! CHECK-NEXT: PolymorphicOpConversion
! CHECK-NEXT: AssumedRankOpConversion
! CHECK-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
! CHECK-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
! CHECK-NEXT: 'fir.global' Pipeline
! CHECK-NEXT: StackReclaim
! CHECK-NEXT: CFGConversion
! CHECK-NEXT: 'func.func' Pipeline
! CHECK-NEXT: StackReclaim
! CHECK-NEXT: CFGConversion
! CHECK-NEXT: 'gpu.module' Pipeline
! CHECK-NEXT: StackReclaim
! CHECK-NEXT: CFGConversion
! CHECK-NEXT: 'omp.declare_reduction' Pipeline
! CHECK-NEXT: StackReclaim
! CHECK-NEXT: CFGConversion

19
flang/test/Driver/mlir-debug-pass-pipeline.f90
View File

@@ -28,13 +28,11 @@ end program
! ALL: Pass statistics report
! ALL: Fortran::lower::VerifierPass
! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
! ALL-NEXT: 'fir.global' Pipeline
! ALL-NEXT: InlineElementals
! ALL-NEXT: 'func.func' Pipeline
! ALL-NEXT: InlineElementals
! ALL-NEXT: 'gpu.module' Pipeline
! ALL-NEXT: InlineElementals
! ALL-NEXT: 'omp.declare_reduction' Pipeline
! ALL-NEXT: InlineElementals
! ALL-NEXT: 'omp.private' Pipeline
@@ -51,14 +49,12 @@ end program
! ALL-NEXT: (S) 0 num-cse'd - Number of operations CSE'd
! ALL-NEXT: (S) 0 num-dce'd - Number of operations DCE'd
! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
! ALL-NEXT: 'fir.global' Pipeline
! ALL-NEXT: CharacterConversion
! ALL-NEXT: 'func.func' Pipeline
! ALL-NEXT: ArrayValueCopy
! ALL-NEXT: CharacterConversion
! ALL-NEXT: 'gpu.module' Pipeline
! ALL-NEXT: CharacterConversion
! ALL-NEXT: 'omp.declare_reduction' Pipeline
! ALL-NEXT: CharacterConversion
! ALL-NEXT: 'omp.private' Pipeline
@@ -82,16 +78,13 @@ end program
! ALL-NEXT: PolymorphicOpConversion
! ALL-NEXT: AssumedRankOpConversion
! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
! ALL-NEXT: 'fir.global' Pipeline
! ALL-NEXT: StackReclaim
! ALL-NEXT: CFGConversion
! ALL-NEXT: 'func.func' Pipeline
! ALL-NEXT: StackReclaim
! ALL-NEXT: CFGConversion
! ALL-NEXT: 'gpu.module' Pipeline
! ALL-NEXT: StackReclaim
! ALL-NEXT: CFGConversion
! ALL-NEXT: 'omp.declare_reduction' Pipeline
! ALL-NEXT: StackReclaim
! ALL-NEXT: CFGConversion
@@ -112,7 +105,11 @@ end program
! ALL-NEXT: 'func.func' Pipeline
! ALL-NEXT: AbstractResultOpt
! ALL-NEXT: 'gpu.module' Pipeline
! ALL-NEXT: AbstractResultOpt
! ALL-NEXT: Pipeline Collection : ['func.func', 'gpu.func']
! ALL-NEXT: 'func.func' Pipeline
! ALL-NEXT: AbstractResultOpt
! ALL-NEXT: 'gpu.func' Pipeline
! ALL-NEXT: AbstractResultOpt
! ALL-NEXT: 'omp.declare_reduction' Pipeline
! ALL-NEXT: AbstractResultOpt
! ALL-NEXT: 'omp.private' Pipeline

24
flang/test/Driver/mlir-pass-pipeline.f90
View File

@@ -16,16 +16,13 @@ end program
! ALL: Fortran::lower::VerifierPass
! O2-NEXT: Canonicalizer
! ALL: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
! ALL: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
! ALL-NEXT:'fir.global' Pipeline
! O2-NEXT: SimplifyHLFIRIntrinsics
! ALL: InlineElementals
! ALL-NEXT:'func.func' Pipeline
! O2-NEXT: SimplifyHLFIRIntrinsics
! ALL: InlineElementals
! ALL-NEXT:'gpu.module' Pipeline
! O2-NEXT: SimplifyHLFIRIntrinsics
! ALL: InlineElementals
! ALL-NEXT:'omp.declare_reduction' Pipeline
! O2-NEXT: SimplifyHLFIRIntrinsics
! ALL: InlineElementals
@@ -36,13 +33,11 @@ end program
! O2-NEXT: CSE
! O2-NEXT: (S) {{.*}} num-cse'd
! O2-NEXT: (S) {{.*}} num-dce'd
! O2-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
! O2-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
! O2-NEXT: 'fir.global' Pipeline
! O2-NEXT: OptimizedBufferization
! O2-NEXT: 'func.func' Pipeline
! O2-NEXT: OptimizedBufferization
! O2-NEXT: 'gpu.module' Pipeline
! O2-NEXT: OptimizedBufferization
! O2-NEXT: 'omp.declare_reduction' Pipeline
! O2-NEXT: OptimizedBufferization
! O2-NEXT: 'omp.private' Pipeline
@@ -59,14 +54,12 @@ end program
! ALL-NEXT: (S) 0 num-cse'd - Number of operations CSE'd
! ALL-NEXT: (S) 0 num-dce'd - Number of operations DCE'd
! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
! ALL-NEXT: 'fir.global' Pipeline
! ALL-NEXT: CharacterConversion
! ALL-NEXT: 'func.func' Pipeline
! ALL-NEXT: ArrayValueCopy
! ALL-NEXT: CharacterConversion
! ALL-NEXT: 'gpu.module' Pipeline
! ALL-NEXT: CharacterConversion
! ALL-NEXT: 'omp.declare_reduction' Pipeline
! ALL-NEXT: CharacterConversion
! ALL-NEXT: 'omp.private' Pipeline
@@ -93,16 +86,13 @@ end program
! ALL-NEXT: AssumedRankOpConversion
! O2-NEXT: AddAliasTags
! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
! ALL-NEXT: 'fir.global' Pipeline
! ALL-NEXT: StackReclaim
! ALL-NEXT: CFGConversion
! ALL-NEXT: 'func.func' Pipeline
! ALL-NEXT: StackReclaim
! ALL-NEXT: CFGConversion
! ALL-NEXT: 'gpu.module' Pipeline
! ALL-NEXT: StackReclaim
! ALL-NEXT: CFGConversion
! ALL-NEXT: 'omp.declare_reduction' Pipeline
! ALL-NEXT: StackReclaim
! ALL-NEXT: CFGConversion
@@ -124,7 +114,11 @@ end program
! ALL-NEXT: 'func.func' Pipeline
! ALL-NEXT: AbstractResultOpt
! ALL-NEXT: 'gpu.module' Pipeline
! ALL-NEXT: AbstractResultOpt
! ALL-NEXT: Pipeline Collection : ['func.func', 'gpu.func']
! ALL-NEXT: 'func.func' Pipeline
! ALL-NEXT: AbstractResultOpt
! ALL-NEXT: 'gpu.func' Pipeline
! ALL-NEXT: AbstractResultOpt
! ALL-NEXT: 'omp.declare_reduction' Pipeline
! ALL-NEXT: AbstractResultOpt
! ALL-NEXT: 'omp.private' Pipeline

24
flang/test/Fir/basic-program.fir
View File

@@ -17,16 +17,13 @@ func.func @_QQmain() {
// PASSES: Pass statistics report
// PASSES: Canonicalizer
// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
// PASSES-NEXT: 'fir.global' Pipeline
// PASSES-NEXT: SimplifyHLFIRIntrinsics
// PASSES-NEXT: InlineElementals
// PASSES-NEXT: 'func.func' Pipeline
// PASSES-NEXT: SimplifyHLFIRIntrinsics
// PASSES-NEXT: InlineElementals
// PASSES-NEXT: 'gpu.module' Pipeline
// PASSES-NEXT: SimplifyHLFIRIntrinsics
// PASSES-NEXT: InlineElementals
// PASSES-NEXT: 'omp.declare_reduction' Pipeline
// PASSES-NEXT: SimplifyHLFIRIntrinsics
// PASSES-NEXT: InlineElementals
@@ -37,13 +34,11 @@ func.func @_QQmain() {
// PASSES-NEXT: CSE
// PASSES-NEXT: (S) 0 num-cse'd - Number of operations CSE'd
// PASSES-NEXT: (S) 0 num-dce'd - Number of operations DCE'd
// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
// PASSES-NEXT: 'fir.global' Pipeline
// PASSES-NEXT: OptimizedBufferization
// PASSES-NEXT: 'func.func' Pipeline
// PASSES-NEXT: OptimizedBufferization
// PASSES-NEXT: 'gpu.module' Pipeline
// PASSES-NEXT: OptimizedBufferization
// PASSES-NEXT: 'omp.declare_reduction' Pipeline
// PASSES-NEXT: OptimizedBufferization
// PASSES-NEXT: 'omp.private' Pipeline
@@ -57,14 +52,12 @@ func.func @_QQmain() {
// PASSES-NEXT: (S) 0 num-cse'd - Number of operations CSE'd
// PASSES-NEXT: (S) 0 num-dce'd - Number of operations DCE'd
// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
// PASSES-NEXT: 'fir.global' Pipeline
// PASSES-NEXT: CharacterConversion
// PASSES-NEXT: 'func.func' Pipeline
// PASSES-NEXT: ArrayValueCopy
// PASSES-NEXT: CharacterConversion
// PASSES-NEXT: 'gpu.module' Pipeline
// PASSES-NEXT: CharacterConversion
// PASSES-NEXT: 'omp.declare_reduction' Pipeline
// PASSES-NEXT: CharacterConversion
// PASSES-NEXT: 'omp.private' Pipeline
@@ -91,16 +84,13 @@ func.func @_QQmain() {
// PASSES-NEXT: AssumedRankOpConversion
// PASSES-NEXT: AddAliasTags
// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private']
// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private']
// PASSES-NEXT: 'fir.global' Pipeline
// PASSES-NEXT: StackReclaim
// PASSES-NEXT: CFGConversion
// PASSES-NEXT: 'func.func' Pipeline
// PASSES-NEXT: StackReclaim
// PASSES-NEXT: CFGConversion
// PASSES-NEXT: 'gpu.module' Pipeline
// PASSES-NEXT: StackReclaim
// PASSES-NEXT: CFGConversion
// PASSES-NEXT: 'omp.declare_reduction' Pipeline
// PASSES-NEXT: StackReclaim
// PASSES-NEXT: CFGConversion
@@ -122,7 +112,11 @@ func.func @_QQmain() {
// PASSES-NEXT: 'func.func' Pipeline
// PASSES-NEXT: AbstractResultOpt
// PASSES-NEXT: 'gpu.module' Pipeline
// PASSES-NEXT: AbstractResultOpt
// PASSES-NEXT: Pipeline Collection : ['func.func', 'gpu.func']
// PASSES-NEXT: 'func.func' Pipeline
// PASSES-NEXT: AbstractResultOpt
// PASSES-NEXT: 'gpu.func' Pipeline
// PASSES-NEXT: AbstractResultOpt
// PASSES-NEXT: 'omp.declare_reduction' Pipeline
// PASSES-NEXT: AbstractResultOpt
// PASSES-NEXT: 'omp.private' Pipeline