1cfd4c180SSlava Zakharin //=== CompilerGeneratedNames.cpp - convert special symbols in global names ===// 2cfd4c180SSlava Zakharin // 3cfd4c180SSlava Zakharin // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4cfd4c180SSlava Zakharin // See https://llvm.org/LICENSE.txt for license information. 5cfd4c180SSlava Zakharin // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6cfd4c180SSlava Zakharin // 7cfd4c180SSlava Zakharin //===----------------------------------------------------------------------===// 8cfd4c180SSlava Zakharin 9cfd4c180SSlava Zakharin #include "flang/Optimizer/Dialect/FIRDialect.h" 10cfd4c180SSlava Zakharin #include "flang/Optimizer/Dialect/FIROps.h" 11cfd4c180SSlava Zakharin #include "flang/Optimizer/Dialect/FIROpsSupport.h" 12cfd4c180SSlava Zakharin #include "flang/Optimizer/Support/InternalNames.h" 13cfd4c180SSlava Zakharin #include "flang/Optimizer/Transforms/Passes.h" 14*e93d2266SValentin Clement (バレンタイン クレメン) #include "mlir/Dialect/GPU/IR/GPUDialect.h" 15cfd4c180SSlava Zakharin #include "mlir/IR/Attributes.h" 16cfd4c180SSlava Zakharin #include "mlir/IR/SymbolTable.h" 17cfd4c180SSlava Zakharin #include "mlir/Pass/Pass.h" 18cfd4c180SSlava Zakharin 19cfd4c180SSlava Zakharin namespace fir { 20cfd4c180SSlava Zakharin #define GEN_PASS_DEF_COMPILERGENERATEDNAMESCONVERSION 21cfd4c180SSlava Zakharin #include "flang/Optimizer/Transforms/Passes.h.inc" 22cfd4c180SSlava Zakharin } // namespace fir 23cfd4c180SSlava Zakharin 24cfd4c180SSlava Zakharin using namespace mlir; 25cfd4c180SSlava Zakharin 26cfd4c180SSlava Zakharin namespace { 27cfd4c180SSlava Zakharin 28cfd4c180SSlava Zakharin class CompilerGeneratedNamesConversionPass 29cfd4c180SSlava Zakharin : public fir::impl::CompilerGeneratedNamesConversionBase< 30cfd4c180SSlava Zakharin CompilerGeneratedNamesConversionPass> { 31cfd4c180SSlava Zakharin public: 32cfd4c180SSlava Zakharin using CompilerGeneratedNamesConversionBase< 33cfd4c180SSlava Zakharin CompilerGeneratedNamesConversionPass>:: 34cfd4c180SSlava Zakharin CompilerGeneratedNamesConversionBase; 35cfd4c180SSlava Zakharin 36cfd4c180SSlava Zakharin mlir::ModuleOp getModule() { return getOperation(); } 37cfd4c180SSlava Zakharin void runOnOperation() override; 38cfd4c180SSlava Zakharin }; 39cfd4c180SSlava Zakharin } // namespace 40cfd4c180SSlava Zakharin 41cfd4c180SSlava Zakharin void CompilerGeneratedNamesConversionPass::runOnOperation() { 42cfd4c180SSlava Zakharin auto op = getOperation(); 43cfd4c180SSlava Zakharin auto *context = &getContext(); 44cfd4c180SSlava Zakharin 45cfd4c180SSlava Zakharin llvm::DenseMap<mlir::StringAttr, mlir::FlatSymbolRefAttr> remappings; 46*e93d2266SValentin Clement (バレンタイン クレメン) 47*e93d2266SValentin Clement (バレンタイン クレメン) auto processOp = [&](mlir::Operation &op) { 48*e93d2266SValentin Clement (バレンタイン クレメン) auto symName = op.getAttrOfType<mlir::StringAttr>( 49cfd4c180SSlava Zakharin mlir::SymbolTable::getSymbolAttrName()); 50cfd4c180SSlava Zakharin auto deconstructedName = fir::NameUniquer::deconstruct(symName); 51cfd4c180SSlava Zakharin if (deconstructedName.first != fir::NameUniquer::NameKind::NOT_UNIQUED && 52cfd4c180SSlava Zakharin !fir::NameUniquer::isExternalFacingUniquedName(deconstructedName)) { 53cfd4c180SSlava Zakharin std::string newName = 54cfd4c180SSlava Zakharin fir::NameUniquer::replaceSpecialSymbols(symName.getValue().str()); 55cfd4c180SSlava Zakharin if (newName != symName) { 56cfd4c180SSlava Zakharin auto newAttr = mlir::StringAttr::get(context, newName); 57*e93d2266SValentin Clement (バレンタイン クレメン) mlir::SymbolTable::setSymbolName(&op, newAttr); 58cfd4c180SSlava Zakharin auto newSymRef = mlir::FlatSymbolRefAttr::get(newAttr); 59cfd4c180SSlava Zakharin remappings.try_emplace(symName, newSymRef); 60cfd4c180SSlava Zakharin } 61cfd4c180SSlava Zakharin } 62*e93d2266SValentin Clement (バレンタイン クレメン) }; 63*e93d2266SValentin Clement (バレンタイン クレメン) for (auto &op : op->getRegion(0).front()) { 64*e93d2266SValentin Clement (バレンタイン クレメン) if (llvm::isa<mlir::func::FuncOp>(op) || llvm::isa<fir::GlobalOp>(op)) 65*e93d2266SValentin Clement (バレンタイン クレメン) processOp(op); 66*e93d2266SValentin Clement (バレンタイン クレメン) else if (auto gpuMod = mlir::dyn_cast<mlir::gpu::GPUModuleOp>(&op)) 67*e93d2266SValentin Clement (バレンタイン クレメン) for (auto &op : gpuMod->getRegion(0).front()) 68*e93d2266SValentin Clement (バレンタイン クレメン) if (llvm::isa<mlir::func::FuncOp>(op) || llvm::isa<fir::GlobalOp>(op) || 69*e93d2266SValentin Clement (バレンタイン クレメン) llvm::isa<mlir::gpu::GPUFuncOp>(op)) 70*e93d2266SValentin Clement (バレンタイン クレメン) processOp(op); 71cfd4c180SSlava Zakharin } 72cfd4c180SSlava Zakharin 73cfd4c180SSlava Zakharin if (remappings.empty()) 74cfd4c180SSlava Zakharin return; 75cfd4c180SSlava Zakharin 76cfd4c180SSlava Zakharin // Update all uses of the functions and globals that have been renamed. 77cfd4c180SSlava Zakharin op.walk([&remappings](mlir::Operation *nestedOp) { 78cfd4c180SSlava Zakharin llvm::SmallVector<std::pair<mlir::StringAttr, mlir::SymbolRefAttr>> updates; 79cfd4c180SSlava Zakharin for (const mlir::NamedAttribute &attr : nestedOp->getAttrDictionary()) 80cfd4c180SSlava Zakharin if (auto symRef = llvm::dyn_cast<mlir::SymbolRefAttr>(attr.getValue())) 81cfd4c180SSlava Zakharin if (auto remap = remappings.find(symRef.getRootReference()); 82cfd4c180SSlava Zakharin remap != remappings.end()) 83cfd4c180SSlava Zakharin updates.emplace_back(std::pair<mlir::StringAttr, mlir::SymbolRefAttr>{ 84cfd4c180SSlava Zakharin attr.getName(), mlir::SymbolRefAttr(remap->second)}); 85cfd4c180SSlava Zakharin for (auto update : updates) 86cfd4c180SSlava Zakharin nestedOp->setAttr(update.first, update.second); 87cfd4c180SSlava Zakharin }); 88cfd4c180SSlava Zakharin } 89