191 lines · cpp
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-exception6//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 20namespace fir {21#define GEN_PASS_DEF_EXTERNALNAMECONVERSION22#include "flang/Optimizer/Transforms/Passes.h.inc"23} // namespace fir24 25using namespace mlir;26 27//===----------------------------------------------------------------------===//28// Helper functions29//===----------------------------------------------------------------------===//30 31/// Mangle the name with gfortran convention.32std::string33mangleExternalName(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/// Process a symbol reference and return the updated symbol reference if45/// needed.46std::optional<mlir::SymbolRefAttr>47processSymbolRef(mlir::SymbolRefAttr symRef, mlir::Operation *nestedOp,48 const llvm::DenseMap<mlir::StringAttr, mlir::FlatSymbolRefAttr>49 &remappings) {50 if (auto remap = remappings.find(symRef.getLeafReference());51 remap != remappings.end()) {52 mlir::SymbolRefAttr symAttr = mlir::FlatSymbolRefAttr(remap->second);53 if (mlir::isa<mlir::gpu::LaunchFuncOp>(nestedOp))54 symAttr = mlir::SymbolRefAttr::get(55 symRef.getRootReference(), {mlir::FlatSymbolRefAttr(remap->second)});56 return symAttr;57 }58 return std::nullopt;59}60 61namespace {62 63class ExternalNameConversionPass64 : public fir::impl::ExternalNameConversionBase<ExternalNameConversionPass> {65public:66 using ExternalNameConversionBase<67 ExternalNameConversionPass>::ExternalNameConversionBase;68 69 mlir::ModuleOp getModule() { return getOperation(); }70 void runOnOperation() override;71};72} // namespace73 74void ExternalNameConversionPass::runOnOperation() {75 auto op = getOperation();76 auto *context = &getContext();77 78 llvm::DenseMap<mlir::StringAttr, mlir::FlatSymbolRefAttr> remappings;79 mlir::SymbolTable symbolTable(op);80 81 auto processFctOrGlobal = [&](mlir::Operation &funcOrGlobal) {82 auto symName = funcOrGlobal.getAttrOfType<mlir::StringAttr>(83 mlir::SymbolTable::getSymbolAttrName());84 auto deconstructedName = fir::NameUniquer::deconstruct(symName);85 if (fir::NameUniquer::isExternalFacingUniquedName(deconstructedName)) {86 // Check if this is a private function that would conflict with a common87 // block and get its mangled name.88 if (auto funcOp = llvm::dyn_cast<mlir::func::FuncOp>(funcOrGlobal)) {89 if (funcOp.isPrivate()) {90 std::string mangledName =91 mangleExternalName(deconstructedName, appendUnderscoreOpt);92 auto mod = funcOp->getParentOfType<mlir::ModuleOp>();93 bool hasConflictingCommonBlock = false;94 95 // Check if any existing global has the same mangled name.96 if (symbolTable.lookup<fir::GlobalOp>(mangledName))97 hasConflictingCommonBlock = true;98 99 // Skip externalization if the function has a conflicting common block100 // and is not directly called (i.e. procedure pointers or type101 // specifications)102 if (hasConflictingCommonBlock) {103 bool isDirectlyCalled = false;104 std::optional<SymbolTable::UseRange> uses =105 funcOp.getSymbolUses(mod);106 if (uses.has_value()) {107 for (auto use : *uses) {108 mlir::Operation *user = use.getUser();109 if (mlir::isa<fir::CallOp>(user) ||110 mlir::isa<mlir::func::CallOp>(user)) {111 isDirectlyCalled = true;112 break;113 }114 }115 }116 if (!isDirectlyCalled)117 return;118 }119 }120 }121 122 auto newName = mangleExternalName(deconstructedName, appendUnderscoreOpt);123 auto newAttr = mlir::StringAttr::get(context, newName);124 mlir::SymbolTable::setSymbolName(&funcOrGlobal, newAttr);125 auto newSymRef = mlir::FlatSymbolRefAttr::get(newAttr);126 remappings.try_emplace(symName, newSymRef);127 if (llvm::isa<mlir::func::FuncOp>(funcOrGlobal))128 funcOrGlobal.setAttr(fir::getInternalFuncNameAttrName(), symName);129 }130 };131 132 auto renameFuncOrGlobalInModule = [&](mlir::Operation *module) {133 for (auto &op : module->getRegion(0).front()) {134 if (mlir::isa<mlir::func::FuncOp, fir::GlobalOp>(op)) {135 processFctOrGlobal(op);136 } else if (auto gpuMod = mlir::dyn_cast<mlir::gpu::GPUModuleOp>(op)) {137 for (auto &gpuOp : gpuMod.getBodyRegion().front())138 if (mlir::isa<mlir::func::FuncOp, fir::GlobalOp,139 mlir::gpu::GPUFuncOp>(gpuOp))140 processFctOrGlobal(gpuOp);141 }142 }143 };144 145 // Update names of external Fortran functions and names of Common Block146 // globals.147 renameFuncOrGlobalInModule(op);148 149 if (remappings.empty())150 return;151 152 // Update all uses of the functions and globals that have been renamed.153 op.walk([&remappings](mlir::Operation *nestedOp) {154 llvm::SmallVector<std::pair<mlir::StringAttr, mlir::SymbolRefAttr>>155 symRefUpdates;156 llvm::SmallVector<std::pair<mlir::StringAttr, mlir::ArrayAttr>>157 arrayUpdates;158 for (const mlir::NamedAttribute &attr : nestedOp->getAttrDictionary())159 if (auto symRef = llvm::dyn_cast<mlir::SymbolRefAttr>(attr.getValue())) {160 if (auto newSymRef = processSymbolRef(symRef, nestedOp, remappings))161 symRefUpdates.emplace_back(162 std::pair<mlir::StringAttr, mlir::SymbolRefAttr>{attr.getName(),163 *newSymRef});164 } else if (auto arrayAttr =165 llvm::dyn_cast<mlir::ArrayAttr>(attr.getValue())) {166 llvm::SmallVector<mlir::Attribute> symbolRefs;167 for (auto element : arrayAttr) {168 if (!element) {169 symbolRefs.push_back(element);170 continue;171 }172 auto symRef = llvm::dyn_cast<mlir::SymbolRefAttr>(element);173 std::optional<mlir::SymbolRefAttr> updatedSymRef;174 if (symRef)175 updatedSymRef = processSymbolRef(symRef, nestedOp, remappings);176 if (!symRef || !updatedSymRef)177 symbolRefs.push_back(element);178 else179 symbolRefs.push_back(*updatedSymRef);180 }181 arrayUpdates.push_back(std::make_pair(182 attr.getName(),183 mlir::ArrayAttr::get(nestedOp->getContext(), symbolRefs)));184 }185 for (auto update : symRefUpdates)186 nestedOp->setAttr(update.first, update.second);187 for (auto update : arrayUpdates)188 nestedOp->setAttr(update.first, update.second);189 });190}191