[flang-commits] [flang] [flang][cuda] Allow AbstractResult to run in gpu.module (PR #118529)
via flang-commits
flang-commits at lists.llvm.org
Tue Dec 3 10:56:00 PST 2024
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-flang-driver
Author: Valentin Clement (バレンタイン クレメン) (clementval)
<details>
<summary>Changes</summary>
in CUDA Fortran, device function are converted to `gpu.func` inside the `gpu.module` operation. Update the AbstractResult pass to be able to run on `func.func` and `gpu.func` operations inside the `gpu.module`.
---
Full diff: https://github.com/llvm/llvm-project/pull/118529.diff
7 Files Affected:
- (modified) flang/include/flang/Optimizer/Passes/Pipelines.h (+1)
- (modified) flang/lib/Optimizer/Passes/Pipelines.cpp (+2-1)
- (modified) flang/lib/Optimizer/Transforms/AbstractResult.cpp (+25-3)
- (modified) flang/test/Driver/bbc-mlir-pass-pipeline.f90 (+7-2)
- (modified) flang/test/Driver/mlir-debug-pass-pipeline.f90 (+13-4)
- (modified) flang/test/Driver/mlir-pass-pipeline.f90 (+17-5)
- (modified) flang/test/Fir/basic-program.fir (+17-5)
``````````diff
diff --git a/flang/include/flang/Optimizer/Passes/Pipelines.h b/flang/include/flang/Optimizer/Passes/Pipelines.h
index 55fafc2e6b36fe..339182605f818f 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 0b7b3bafde008c..0743fb60aa847a 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<mlir::func::FuncOp, mlir::omp::DeclareReductionOp,
- mlir::omp::PrivateClauseOp, fir::GlobalOp>(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 e64280508755a4..2eca349110f3af 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 <typename OpTy>
+ 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<mlir::Type, bool>(type)
.Case([](fir::BoxProcType boxProc) {
@@ -448,6 +462,14 @@ class AbstractResultOpt
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);
});
// 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 5520d750e2ce1c..1f09e7ad4c2f5a 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 ab5ddedf5fc180..4326953421e4bd 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 33c8183b27aef3..6ffdbb0234e856 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 ad5201af8311ff..50b91ce340b3a6 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
``````````
</details>
https://github.com/llvm/llvm-project/pull/118529
More information about the flang-commits
mailing list