diff --git a/flang/include/flang/Optimizer/Passes/Pipelines.h b/flang/include/flang/Optimizer/Passes/Pipelines.h index 55fafc2e6b36f..339182605f818 100644 --- a/flang/include/flang/Optimizer/Passes/Pipelines.h +++ b/flang/include/flang/Optimizer/Passes/Pipelines.h @@ -20,6 +20,7 @@ #include "flang/Tools/CrossToolHelpers.h" #include "mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h" #include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h" +#include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/LLVMIR/LLVMAttrs.h" #include "mlir/Pass/PassManager.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" diff --git a/flang/lib/Optimizer/Passes/Pipelines.cpp b/flang/lib/Optimizer/Passes/Pipelines.cpp index 0b7b3bafde008..0743fb60aa847 100644 --- a/flang/lib/Optimizer/Passes/Pipelines.cpp +++ b/flang/lib/Optimizer/Passes/Pipelines.cpp @@ -16,7 +16,8 @@ namespace fir { void addNestedPassToAllTopLevelOperations(mlir::PassManager &pm, PassConstructor ctor) { addNestedPassToOps(pm, ctor); + mlir::omp::PrivateClauseOp, fir::GlobalOp, + mlir::gpu::GPUModuleOp>(pm, ctor); } void addNestedPassToAllTopLevelOperationsConditionally( diff --git a/flang/lib/Optimizer/Transforms/AbstractResult.cpp b/flang/lib/Optimizer/Transforms/AbstractResult.cpp index e64280508755a..2eca349110f3a 100644 --- a/flang/lib/Optimizer/Transforms/AbstractResult.cpp +++ b/flang/lib/Optimizer/Transforms/AbstractResult.cpp @@ -14,6 +14,7 @@ #include "flang/Optimizer/Dialect/Support/FIRContext.h" #include "flang/Optimizer/Transforms/Passes.h" #include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/IR/Diagnostics.h" #include "mlir/Pass/Pass.h" #include "mlir/Pass/PassManager.h" @@ -331,9 +332,10 @@ class AbstractResultOpt using fir::impl::AbstractResultOptBase< AbstractResultOpt>::AbstractResultOptBase; - void runOnSpecificOperation(mlir::func::FuncOp func, bool shouldBoxResult, - mlir::RewritePatternSet &patterns, - mlir::ConversionTarget &target) { + template + void runOnFunctionLikeOperation(OpTy func, bool shouldBoxResult, + mlir::RewritePatternSet &patterns, + mlir::ConversionTarget &target) { auto loc = func.getLoc(); auto *context = &getContext(); // Convert function type itself if it has an abstract result. @@ -384,6 +386,18 @@ class AbstractResultOpt } } + void runOnSpecificOperation(mlir::func::FuncOp func, bool shouldBoxResult, + mlir::RewritePatternSet &patterns, + mlir::ConversionTarget &target) { + runOnFunctionLikeOperation(func, shouldBoxResult, patterns, target); + } + + void runOnSpecificOperation(mlir::gpu::GPUFuncOp func, bool shouldBoxResult, + mlir::RewritePatternSet &patterns, + mlir::ConversionTarget &target) { + runOnFunctionLikeOperation(func, shouldBoxResult, patterns, target); + } + inline static bool containsFunctionTypeWithAbstractResult(mlir::Type type) { return mlir::TypeSwitch(type) .Case([](fir::BoxProcType boxProc) { @@ -448,6 +462,14 @@ class AbstractResultOpt mlir::TypeSwitch(op) .Case([&](auto op) { runOnSpecificOperation(op, shouldBoxResult, patterns, target); + }) + .Case([&](auto op) { + auto gpuMod = mlir::dyn_cast(*op); + for (auto funcOp : gpuMod.template getOps()) + runOnSpecificOperation(funcOp, shouldBoxResult, patterns, target); + for (auto gpuFuncOp : gpuMod.template getOps()) + runOnSpecificOperation(gpuFuncOp, shouldBoxResult, patterns, + target); }); // Convert the calls and, if needed, the ReturnOp in the function body. diff --git a/flang/test/Driver/bbc-mlir-pass-pipeline.f90 b/flang/test/Driver/bbc-mlir-pass-pipeline.f90 index 5520d750e2ce1..1f09e7ad4c2f5 100644 --- a/flang/test/Driver/bbc-mlir-pass-pipeline.f90 +++ b/flang/test/Driver/bbc-mlir-pass-pipeline.f90 @@ -17,12 +17,14 @@ ! 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', 'omp.declare_reduction', 'omp.private'] +! CHECK-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 @@ -48,13 +50,16 @@ ! CHECK-NEXT: PolymorphicOpConversion ! CHECK-NEXT: AssumedRankOpConversion -! CHECK-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private'] +! CHECK-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 diff --git a/flang/test/Driver/mlir-debug-pass-pipeline.f90 b/flang/test/Driver/mlir-debug-pass-pipeline.f90 index ab5ddedf5fc18..4326953421e4b 100644 --- a/flang/test/Driver/mlir-debug-pass-pipeline.f90 +++ b/flang/test/Driver/mlir-debug-pass-pipeline.f90 @@ -28,11 +28,13 @@ ! ALL: Pass statistics report ! ALL: Fortran::lower::VerifierPass -! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private'] +! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 @@ -49,12 +51,14 @@ ! 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', 'omp.declare_reduction', 'omp.private'] +! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 @@ -78,13 +82,16 @@ ! ALL-NEXT: PolymorphicOpConversion ! ALL-NEXT: AssumedRankOpConversion -! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private'] +! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 @@ -99,11 +106,13 @@ ! ALL-NEXT: (S) 0 num-dce'd - Number of operations DCE'd ! ALL-NEXT: BoxedProcedurePass -! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private'] +! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private'] ! ALL-NEXT: 'fir.global' Pipeline ! ALL-NEXT: AbstractResultOpt ! ALL-NEXT: 'func.func' Pipeline ! ALL-NEXT: AbstractResultOpt +! ALL-NEXT: 'gpu.module' Pipeline +! ALL-NEXT: AbstractResultOpt ! ALL-NEXT: 'omp.declare_reduction' Pipeline ! ALL-NEXT: AbstractResultOpt ! ALL-NEXT: 'omp.private' Pipeline diff --git a/flang/test/Driver/mlir-pass-pipeline.f90 b/flang/test/Driver/mlir-pass-pipeline.f90 index 33c8183b27aef..6ffdbb0234e85 100644 --- a/flang/test/Driver/mlir-pass-pipeline.f90 +++ b/flang/test/Driver/mlir-pass-pipeline.f90 @@ -16,13 +16,16 @@ ! ALL: Fortran::lower::VerifierPass ! O2-NEXT: Canonicalizer -! ALL: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private'] +! ALL: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 @@ -33,11 +36,13 @@ ! O2-NEXT: CSE ! O2-NEXT: (S) {{.*}} num-cse'd ! O2-NEXT: (S) {{.*}} num-dce'd -! O2-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private'] +! O2-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 @@ -54,12 +59,14 @@ ! 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', 'omp.declare_reduction', 'omp.private'] +! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 @@ -86,13 +93,16 @@ ! ALL-NEXT: AssumedRankOpConversion ! O2-NEXT: AddAliasTags -! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private'] +! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 @@ -108,11 +118,13 @@ ! ALL-NEXT: (S) 0 num-dce'd - Number of operations DCE'd ! ALL-NEXT: BoxedProcedurePass -! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private'] +! ALL-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private'] ! ALL-NEXT: 'fir.global' Pipeline ! ALL-NEXT: AbstractResultOpt ! ALL-NEXT: 'func.func' Pipeline ! ALL-NEXT: AbstractResultOpt +! ALL-NEXT: 'gpu.module' Pipeline +! ALL-NEXT: AbstractResultOpt ! ALL-NEXT: 'omp.declare_reduction' Pipeline ! ALL-NEXT: AbstractResultOpt ! ALL-NEXT: 'omp.private' Pipeline diff --git a/flang/test/Fir/basic-program.fir b/flang/test/Fir/basic-program.fir index ad5201af8311f..50b91ce340b3a 100644 --- a/flang/test/Fir/basic-program.fir +++ b/flang/test/Fir/basic-program.fir @@ -17,13 +17,16 @@ func.func @_QQmain() { // PASSES: Pass statistics report // PASSES: Canonicalizer -// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private'] +// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 @@ -34,11 +37,13 @@ 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', 'omp.declare_reduction', 'omp.private'] +// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 @@ -52,12 +57,14 @@ 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', 'omp.declare_reduction', 'omp.private'] +// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 @@ -84,13 +91,16 @@ func.func @_QQmain() { // PASSES-NEXT: AssumedRankOpConversion // PASSES-NEXT: AddAliasTags -// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private'] +// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', '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 @@ -106,11 +116,13 @@ func.func @_QQmain() { // PASSES-NEXT: (S) 0 num-dce'd - Number of operations DCE'd // PASSES-NEXT: BoxedProcedurePass -// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'omp.declare_reduction', 'omp.private'] +// PASSES-NEXT: Pipeline Collection : ['fir.global', 'func.func', 'gpu.module', 'omp.declare_reduction', 'omp.private'] // PASSES-NEXT: 'fir.global' Pipeline // PASSES-NEXT: AbstractResultOpt // PASSES-NEXT: 'func.func' Pipeline // PASSES-NEXT: AbstractResultOpt +// PASSES-NEXT: 'gpu.module' Pipeline +// PASSES-NEXT: AbstractResultOpt // PASSES-NEXT: 'omp.declare_reduction' Pipeline // PASSES-NEXT: AbstractResultOpt // PASSES-NEXT: 'omp.private' Pipeline