diff options
| author | Valentin Clement (バレンタイン クレメン) <clementval@gmail.com> | 2024-12-04 10:00:42 -0800 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2024-12-04 10:00:42 -0800 |
| commit | cd92c6a89541cbbb67b39142d93a76caae0f79bf (patch) | |
| tree | 7a8796b5dfedb9829d7de4635a41493596139ee8 /flang/lib/Optimizer/CodeGen/TargetRewrite.cpp | |
| parent | d057b53a7db43f33f4a9fd832115e613ebe0a67b (diff) | |
| download | llvm-cd92c6a89541cbbb67b39142d93a76caae0f79bf.zip llvm-cd92c6a89541cbbb67b39142d93a76caae0f79bf.tar.gz llvm-cd92c6a89541cbbb67b39142d93a76caae0f79bf.tar.bz2 | |
[flang][cuda] Run target rewrite in gpu.module (#118592)
Apply signature conversion for `func.func` in the gpu.module. More work
will need to be done for gpu.func op and implement the NVVM ABI for
conversion in the gpu module.
Diffstat (limited to 'flang/lib/Optimizer/CodeGen/TargetRewrite.cpp')
| -rw-r--r-- | flang/lib/Optimizer/CodeGen/TargetRewrite.cpp | 6 |
1 files changed, 6 insertions, 0 deletions
diff --git a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp index ae6e7ce..1b86d52 100644 --- a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp +++ b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp @@ -27,6 +27,7 @@ #include "flang/Optimizer/Dialect/Support/FIRContext.h" #include "flang/Optimizer/Support/DataLayout.h" #include "mlir/Dialect/DLTI/DLTI.h" +#include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Transforms/DialectConversion.h" #include "llvm/ADT/STLExtras.h" @@ -720,6 +721,11 @@ public: convertSignature(fn); } + + for (auto gpuMod : mod.getOps<mlir::gpu::GPUModuleOp>()) + for (auto fn : gpuMod.getOps<mlir::func::FuncOp>()) + convertSignature(fn); + return mlir::success(); } |