1 | //===- ExternalNameConversion.cpp -- convert name with external convention ===// |
2 | // |
3 | // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. |
4 | // See https://llvm.org/LICENSE.txt for license information. |
5 | // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception |
6 | // |
7 | //===----------------------------------------------------------------------===// |
8 | |
9 | #include "flang/Optimizer/Dialect/FIRDialect.h" |
10 | #include "flang/Optimizer/Dialect/FIROps.h" |
11 | #include "flang/Optimizer/Dialect/FIROpsSupport.h" |
12 | #include "flang/Optimizer/Support/InternalNames.h" |
13 | #include "flang/Optimizer/Transforms/Passes.h" |
14 | #include "flang/Support/Fortran.h" |
15 | #include "mlir/Dialect/GPU/IR/GPUDialect.h" |
16 | #include "mlir/IR/Attributes.h" |
17 | #include "mlir/IR/SymbolTable.h" |
18 | #include "mlir/Pass/Pass.h" |
19 | |
20 | namespace fir { |
21 | #define GEN_PASS_DEF_EXTERNALNAMECONVERSION |
22 | #include "flang/Optimizer/Transforms/Passes.h.inc" |
23 | } // namespace fir |
24 | |
25 | using namespace mlir; |
26 | |
27 | //===----------------------------------------------------------------------===// |
28 | // Helper functions |
29 | //===----------------------------------------------------------------------===// |
30 | |
31 | /// Mangle the name with gfortran convention. |
32 | std::string |
33 | mangleExternalName(const std::pair<fir::NameUniquer::NameKind, |
34 | fir::NameUniquer::DeconstructedName> |
35 | result, |
36 | bool appendUnderscore) { |
37 | if (result.first == fir::NameUniquer::NameKind::COMMON && |
38 | result.second.name.empty()) |
39 | return Fortran::common::blankCommonObjectName; |
40 | return Fortran::common::GetExternalAssemblyName(result.second.name, |
41 | appendUnderscore); |
42 | } |
43 | |
44 | namespace { |
45 | |
46 | class ExternalNameConversionPass |
47 | : public fir::impl::ExternalNameConversionBase<ExternalNameConversionPass> { |
48 | public: |
49 | using ExternalNameConversionBase< |
50 | ExternalNameConversionPass>::ExternalNameConversionBase; |
51 | |
52 | mlir::ModuleOp getModule() { return getOperation(); } |
53 | void runOnOperation() override; |
54 | }; |
55 | } // namespace |
56 | |
57 | void ExternalNameConversionPass::runOnOperation() { |
58 | auto op = getOperation(); |
59 | auto *context = &getContext(); |
60 | |
61 | llvm::DenseMap<mlir::StringAttr, mlir::FlatSymbolRefAttr> remappings; |
62 | |
63 | auto processFctOrGlobal = [&](mlir::Operation &funcOrGlobal) { |
64 | auto symName = funcOrGlobal.getAttrOfType<mlir::StringAttr>( |
65 | mlir::SymbolTable::getSymbolAttrName()); |
66 | auto deconstructedName = fir::NameUniquer::deconstruct(symName); |
67 | if (fir::NameUniquer::isExternalFacingUniquedName(deconstructedName)) { |
68 | auto newName = mangleExternalName(deconstructedName, appendUnderscoreOpt); |
69 | auto newAttr = mlir::StringAttr::get(context, newName); |
70 | mlir::SymbolTable::setSymbolName(&funcOrGlobal, newAttr); |
71 | auto newSymRef = mlir::FlatSymbolRefAttr::get(newAttr); |
72 | remappings.try_emplace(symName, newSymRef); |
73 | if (llvm::isa<mlir::func::FuncOp>(funcOrGlobal)) |
74 | funcOrGlobal.setAttr(fir::getInternalFuncNameAttrName(), symName); |
75 | } |
76 | }; |
77 | |
78 | auto renameFuncOrGlobalInModule = [&](mlir::Operation *module) { |
79 | for (auto &op : module->getRegion(0).front()) { |
80 | if (mlir::isa<mlir::func::FuncOp, fir::GlobalOp>(op)) { |
81 | processFctOrGlobal(op); |
82 | } else if (auto gpuMod = mlir::dyn_cast<mlir::gpu::GPUModuleOp>(op)) { |
83 | for (auto &gpuOp : gpuMod.getBodyRegion().front()) |
84 | if (mlir::isa<mlir::func::FuncOp, fir::GlobalOp, |
85 | mlir::gpu::GPUFuncOp>(gpuOp)) |
86 | processFctOrGlobal(gpuOp); |
87 | } |
88 | } |
89 | }; |
90 | |
91 | // Update names of external Fortran functions and names of Common Block |
92 | // globals. |
93 | renameFuncOrGlobalInModule(op); |
94 | |
95 | if (remappings.empty()) |
96 | return; |
97 | |
98 | // Update all uses of the functions and globals that have been renamed. |
99 | op.walk([&remappings](mlir::Operation *nestedOp) { |
100 | llvm::SmallVector<std::pair<mlir::StringAttr, mlir::SymbolRefAttr>> updates; |
101 | for (const mlir::NamedAttribute &attr : nestedOp->getAttrDictionary()) |
102 | if (auto symRef = llvm::dyn_cast<mlir::SymbolRefAttr>(attr.getValue())) { |
103 | if (auto remap = remappings.find(symRef.getLeafReference()); |
104 | remap != remappings.end()) { |
105 | mlir::SymbolRefAttr symAttr = mlir::FlatSymbolRefAttr(remap->second); |
106 | if (mlir::isa<mlir::gpu::LaunchFuncOp>(nestedOp)) |
107 | symAttr = mlir::SymbolRefAttr::get( |
108 | symRef.getRootReference(), |
109 | {mlir::FlatSymbolRefAttr(remap->second)}); |
110 | updates.emplace_back(std::pair<mlir::StringAttr, mlir::SymbolRefAttr>{ |
111 | attr.getName(), symAttr}); |
112 | } |
113 | } |
114 | for (auto update : updates) |
115 | nestedOp->setAttr(update.first, update.second); |
116 | }); |
117 | } |
118 | |