[flang-commits] [flang] e93d226 - [flang][cuda] Update CompilerGeneratedNames pass to work on gpu module (#120660)
via flang-commits
flang-commits at lists.llvm.org
Thu Dec 19 19:07:04 PST 2024
Author: Valentin Clement (バレンタイン クレメン)
Date: 2024-12-19T19:07:00-08:00
New Revision: e93d226664d7012d1bb017f0cda24ad1b75f37fc
URL: https://github.com/llvm/llvm-project/commit/e93d226664d7012d1bb017f0cda24ad1b75f37fc
DIFF: https://github.com/llvm/llvm-project/commit/e93d226664d7012d1bb017f0cda24ad1b75f37fc.diff
LOG: [flang][cuda] Update CompilerGeneratedNames pass to work on gpu module (#120660)
- Update `CompilerGeneratedNames` so it can perform renaming in
gpu.module
- Update Codegen so it look in the correct module for the type
descriptor.
Added:
flang/test/Fir/CUDA/cuda-compiler-generated-names.mlir
Modified:
flang/lib/Optimizer/CodeGen/CodeGen.cpp
flang/lib/Optimizer/Transforms/CompilerGeneratedNames.cpp
Removed:
################################################################################
diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index b3f864f4f2d4fb..aaf97d46d83d4f 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -1341,10 +1341,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())
@@ -1369,7 +1369,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,
@@ -1508,10 +1509,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};
}
@@ -1554,11 +1561,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