[flang-commits] [flang] [flang][cuda] Update CompilerGeneratedNames pass to work on gpu module (PR #120660)
Valentin Clement バレンタイン クレメン via flang-commits
flang-commits at lists.llvm.org
Thu Dec 19 16:47:23 PST 2024
https://github.com/clementval created https://github.com/llvm/llvm-project/pull/120660
- Update `CompilerGeneratedNames` so it can perform renaming in gpu.module
- Update Codegen so it look in the correct module for the type descriptor.
>From 537359ee4484e83af02210b80ee39cbfefd2d0bd Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 19 Dec 2024 16:42:41 -0800
Subject: [PATCH] [flang][cuda] Update CompilerGeneratedNames pass to work on
gpu module
---
flang/lib/Optimizer/CodeGen/CodeGen.cpp | 41 ++++++++++++-------
.../Transforms/CompilerGeneratedNames.cpp | 40 ++++++++++--------
.../CUDA/cuda-compiler-generated-names.mlir | 17 ++++++++
3 files changed, 68 insertions(+), 30 deletions(-)
create mode 100644 flang/test/Fir/CUDA/cuda-compiler-generated-names.mlir
diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index 082f2b15512b8b..6958356add8137 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -1247,10 +1247,10 @@ struct EmboxCommonConversion : public fir::FIROpConversion<OP> {
/// Get the address of the type descriptor global variable that was created by
/// lowering for derived type \p recType.
- mlir::Value getTypeDescriptor(mlir::ModuleOp mod,
- mlir::ConversionPatternRewriter &rewriter,
- mlir::Location loc,
- fir::RecordType recType) const {
+ template <typename ModOpTy>
+ mlir::Value
+ getTypeDescriptor(ModOpTy mod, mlir::ConversionPatternRewriter &rewriter,
+ mlir::Location loc, fir::RecordType recType) const {
std::string name =
this->options.typeDescriptorsRenamedForAssembly
? fir::NameUniquer::getTypeDescriptorAssemblyName(recType.getName())
@@ -1275,7 +1275,8 @@ struct EmboxCommonConversion : public fir::FIROpConversion<OP> {
return rewriter.create<mlir::LLVM::ZeroOp>(loc, llvmPtrTy);
}
- mlir::Value populateDescriptor(mlir::Location loc, mlir::ModuleOp mod,
+ template <typename ModOpTy>
+ mlir::Value populateDescriptor(mlir::Location loc, ModOpTy mod,
fir::BaseBoxType boxTy, mlir::Type inputType,
mlir::ConversionPatternRewriter &rewriter,
unsigned rank, mlir::Value eleSize,
@@ -1414,10 +1415,16 @@ struct EmboxCommonConversion : public fir::FIROpConversion<OP> {
extraField =
this->getExtraFromBox(loc, sourceBoxTyPair, sourceBox, rewriter);
}
- auto mod = box->template getParentOfType<mlir::ModuleOp>();
- mlir::Value descriptor =
- populateDescriptor(loc, mod, boxTy, inputType, rewriter, rank, eleSize,
- cfiTy, typeDesc, allocatorIdx, extraField);
+
+ mlir::Value descriptor;
+ if (auto gpuMod = box->template getParentOfType<mlir::gpu::GPUModuleOp>())
+ descriptor = populateDescriptor(loc, gpuMod, boxTy, inputType, rewriter,
+ rank, eleSize, cfiTy, typeDesc,
+ allocatorIdx, extraField);
+ else if (auto mod = box->template getParentOfType<mlir::ModuleOp>())
+ descriptor = populateDescriptor(loc, mod, boxTy, inputType, rewriter,
+ rank, eleSize, cfiTy, typeDesc,
+ allocatorIdx, extraField);
return {boxTy, descriptor, eleSize};
}
@@ -1460,11 +1467,17 @@ struct EmboxCommonConversion : public fir::FIROpConversion<OP> {
mlir::Value extraField =
this->getExtraFromBox(loc, inputBoxTyPair, loweredBox, rewriter);
- auto mod = box->template getParentOfType<mlir::ModuleOp>();
- mlir::Value descriptor =
- populateDescriptor(loc, mod, boxTy, box.getBox().getType(), rewriter,
- rank, eleSize, cfiTy, typeDesc,
- /*allocatorIdx=*/kDefaultAllocator, extraField);
+ mlir::Value descriptor;
+ if (auto gpuMod = box->template getParentOfType<mlir::gpu::GPUModuleOp>())
+ descriptor =
+ populateDescriptor(loc, gpuMod, boxTy, box.getBox().getType(),
+ rewriter, rank, eleSize, cfiTy, typeDesc,
+ /*allocatorIdx=*/kDefaultAllocator, extraField);
+ else if (auto mod = box->template getParentOfType<mlir::ModuleOp>())
+ descriptor =
+ populateDescriptor(loc, mod, boxTy, box.getBox().getType(), rewriter,
+ rank, eleSize, cfiTy, typeDesc,
+ /*allocatorIdx=*/kDefaultAllocator, extraField);
return {boxTy, descriptor, eleSize};
}
diff --git a/flang/lib/Optimizer/Transforms/CompilerGeneratedNames.cpp b/flang/lib/Optimizer/Transforms/CompilerGeneratedNames.cpp
index 7f2cc41275e593..f92c60908b1496 100644
--- a/flang/lib/Optimizer/Transforms/CompilerGeneratedNames.cpp
+++ b/flang/lib/Optimizer/Transforms/CompilerGeneratedNames.cpp
@@ -11,6 +11,7 @@
#include "flang/Optimizer/Dialect/FIROpsSupport.h"
#include "flang/Optimizer/Support/InternalNames.h"
#include "flang/Optimizer/Transforms/Passes.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/SymbolTable.h"
#include "mlir/Pass/Pass.h"
@@ -42,24 +43,31 @@ void CompilerGeneratedNamesConversionPass::runOnOperation() {
auto *context = &getContext();
llvm::DenseMap<mlir::StringAttr, mlir::FlatSymbolRefAttr> remappings;
- for (auto &funcOrGlobal : op->getRegion(0).front()) {
- if (llvm::isa<mlir::func::FuncOp>(funcOrGlobal) ||
- llvm::isa<fir::GlobalOp>(funcOrGlobal)) {
- auto symName = funcOrGlobal.getAttrOfType<mlir::StringAttr>(
- mlir::SymbolTable::getSymbolAttrName());
- auto deconstructedName = fir::NameUniquer::deconstruct(symName);
- if (deconstructedName.first != fir::NameUniquer::NameKind::NOT_UNIQUED &&
- !fir::NameUniquer::isExternalFacingUniquedName(deconstructedName)) {
- std::string newName =
- fir::NameUniquer::replaceSpecialSymbols(symName.getValue().str());
- if (newName != symName) {
- auto newAttr = mlir::StringAttr::get(context, newName);
- mlir::SymbolTable::setSymbolName(&funcOrGlobal, newAttr);
- auto newSymRef = mlir::FlatSymbolRefAttr::get(newAttr);
- remappings.try_emplace(symName, newSymRef);
- }
+
+ auto processOp = [&](mlir::Operation &op) {
+ auto symName = op.getAttrOfType<mlir::StringAttr>(
+ mlir::SymbolTable::getSymbolAttrName());
+ auto deconstructedName = fir::NameUniquer::deconstruct(symName);
+ if (deconstructedName.first != fir::NameUniquer::NameKind::NOT_UNIQUED &&
+ !fir::NameUniquer::isExternalFacingUniquedName(deconstructedName)) {
+ std::string newName =
+ fir::NameUniquer::replaceSpecialSymbols(symName.getValue().str());
+ if (newName != symName) {
+ auto newAttr = mlir::StringAttr::get(context, newName);
+ mlir::SymbolTable::setSymbolName(&op, newAttr);
+ auto newSymRef = mlir::FlatSymbolRefAttr::get(newAttr);
+ remappings.try_emplace(symName, newSymRef);
}
}
+ };
+ for (auto &op : op->getRegion(0).front()) {
+ if (llvm::isa<mlir::func::FuncOp>(op) || llvm::isa<fir::GlobalOp>(op))
+ processOp(op);
+ else if (auto gpuMod = mlir::dyn_cast<mlir::gpu::GPUModuleOp>(&op))
+ for (auto &op : gpuMod->getRegion(0).front())
+ if (llvm::isa<mlir::func::FuncOp>(op) || llvm::isa<fir::GlobalOp>(op) ||
+ llvm::isa<mlir::gpu::GPUFuncOp>(op))
+ processOp(op);
}
if (remappings.empty())
diff --git a/flang/test/Fir/CUDA/cuda-compiler-generated-names.mlir b/flang/test/Fir/CUDA/cuda-compiler-generated-names.mlir
new file mode 100644
index 00000000000000..4507e444d1b510
--- /dev/null
+++ b/flang/test/Fir/CUDA/cuda-compiler-generated-names.mlir
@@ -0,0 +1,17 @@
+// RUN: fir-opt --split-input-file --compiler-generated-names --fir-to-llvm-ir="target=x86_64-unknown-linux-gnu type-descriptors-renamed-for-assembly=true" %s | FileCheck %s
+
+module @mod1 attributes {gpu.container} {
+ gpu.module @gpu1 {
+ fir.global linkonce @_QMtest_dinitE.dt.tseq constant : i8
+
+ func.func @embox1(%arg0: !fir.ref<!fir.type<_QMtest_dinitTtseq{i:i32}>>) {
+ %0 = fir.embox %arg0() : (!fir.ref<!fir.type<_QMtest_dinitTtseq{i:i32}>>) -> !fir.box<!fir.type<_QMtest_dinitTtseq{i:i32}>>
+ return
+ }
+ }
+}
+
+// CHECK-LABEL: gpu.module @gpu1
+// CHECK: llvm.mlir.global linkonce constant @_QMtest_dinitEXdtXtseq
+// CHECK: llvm.mlir.addressof @_QMtest_dinitEXdtXtseq : !llvm.ptr
+
More information about the flang-commits
mailing list