blob: f92c60908b149622a1b3b55aaf3ad1fd84ccff6c (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
|
//=== CompilerGeneratedNames.cpp - convert special symbols in global names ===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
#include "flang/Optimizer/Dialect/FIRDialect.h"
#include "flang/Optimizer/Dialect/FIROps.h"
#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"
namespace fir {
#define GEN_PASS_DEF_COMPILERGENERATEDNAMESCONVERSION
#include "flang/Optimizer/Transforms/Passes.h.inc"
} // namespace fir
using namespace mlir;
namespace {
class CompilerGeneratedNamesConversionPass
: public fir::impl::CompilerGeneratedNamesConversionBase<
CompilerGeneratedNamesConversionPass> {
public:
using CompilerGeneratedNamesConversionBase<
CompilerGeneratedNamesConversionPass>::
CompilerGeneratedNamesConversionBase;
mlir::ModuleOp getModule() { return getOperation(); }
void runOnOperation() override;
};
} // namespace
void CompilerGeneratedNamesConversionPass::runOnOperation() {
auto op = getOperation();
auto *context = &getContext();
llvm::DenseMap<mlir::StringAttr, mlir::FlatSymbolRefAttr> remappings;
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())
return;
// Update all uses of the functions and globals that have been renamed.
op.walk([&remappings](mlir::Operation *nestedOp) {
llvm::SmallVector<std::pair<mlir::StringAttr, mlir::SymbolRefAttr>> updates;
for (const mlir::NamedAttribute &attr : nestedOp->getAttrDictionary())
if (auto symRef = llvm::dyn_cast<mlir::SymbolRefAttr>(attr.getValue()))
if (auto remap = remappings.find(symRef.getRootReference());
remap != remappings.end())
updates.emplace_back(std::pair<mlir::StringAttr, mlir::SymbolRefAttr>{
attr.getName(), mlir::SymbolRefAttr(remap->second)});
for (auto update : updates)
nestedOp->setAttr(update.first, update.second);
});
}
|
