[flang][cuda] Update CompilerGeneratedNames pass to work on gpu module #120660
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
llvmbot
added
flang
|
@llvm/pr-subscribers-flang-fir-hlfir @llvm/pr-subscribers-flang-codegen Author: Valentin Clement (バレンタイン クレメン) (clementval) Changes
Full diff: https://github.com/llvm/llvm-project/pull/120660.diff 3 Files Affected:
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
+
|
|
I can split that in two patches if necessary
|
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
CompilerGeneratedNamesso it can perform renaming in gpu.module