1401 lines · cpp
1//===-- TargetRewrite.cpp -------------------------------------------------===//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// Target rewrite: rewriting of ops to make target-specific lowerings manifest.10// LLVM expects different lowering idioms to be used for distinct target11// triples. These distinctions are handled by this pass.12//13// Coding style: https://mlir.llvm.org/getting_started/DeveloperGuide/14//15//===----------------------------------------------------------------------===//16 17#include "flang/Optimizer/CodeGen/CodeGen.h"18 19#include "flang/Optimizer/Builder/Character.h"20#include "flang/Optimizer/Builder/FIRBuilder.h"21#include "flang/Optimizer/Builder/Todo.h"22#include "flang/Optimizer/CodeGen/Target.h"23#include "flang/Optimizer/Dialect/FIRDialect.h"24#include "flang/Optimizer/Dialect/FIROps.h"25#include "flang/Optimizer/Dialect/FIROpsSupport.h"26#include "flang/Optimizer/Dialect/FIRType.h"27#include "flang/Optimizer/Dialect/Support/FIRContext.h"28#include "flang/Optimizer/Support/DataLayout.h"29#include "mlir/Dialect/DLTI/DLTI.h"30#include "mlir/Dialect/GPU/IR/GPUDialect.h"31#include "mlir/Dialect/LLVMIR/LLVMDialect.h"32#include "mlir/Transforms/DialectConversion.h"33#include "llvm/ADT/STLExtras.h"34#include "llvm/ADT/TypeSwitch.h"35#include "llvm/Support/Debug.h"36#include <optional>37 38namespace fir {39#define GEN_PASS_DEF_TARGETREWRITEPASS40#include "flang/Optimizer/CodeGen/CGPasses.h.inc"41} // namespace fir42 43#define DEBUG_TYPE "flang-target-rewrite"44 45namespace {46 47/// Fixups for updating a FuncOp's arguments and return values.48struct FixupTy {49 enum class Codes {50 ArgumentAsLoad,51 ArgumentType,52 CharPair,53 ReturnAsStore,54 ReturnType,55 Split,56 Trailing,57 TrailingCharProc58 };59 60 FixupTy(Codes code, std::size_t index, std::size_t second = 0)61 : code{code}, index{index}, second{second} {}62 FixupTy(Codes code, std::size_t index,63 std::function<void(mlir::func::FuncOp)> &&finalizer)64 : code{code}, index{index}, finalizer{finalizer} {}65 FixupTy(Codes code, std::size_t index,66 std::function<void(mlir::gpu::GPUFuncOp)> &&finalizer)67 : code{code}, index{index}, gpuFinalizer{finalizer} {}68 FixupTy(Codes code, std::size_t index, std::size_t second,69 std::function<void(mlir::func::FuncOp)> &&finalizer)70 : code{code}, index{index}, second{second}, finalizer{finalizer} {}71 FixupTy(Codes code, std::size_t index, std::size_t second,72 std::function<void(mlir::gpu::GPUFuncOp)> &&finalizer)73 : code{code}, index{index}, second{second}, gpuFinalizer{finalizer} {}74 75 Codes code;76 std::size_t index;77 std::size_t second{};78 std::optional<std::function<void(mlir::func::FuncOp)>> finalizer{};79 std::optional<std::function<void(mlir::gpu::GPUFuncOp)>> gpuFinalizer{};80}; // namespace81 82/// Target-specific rewriting of the FIR. This is a prerequisite pass to code83/// generation that traverses the FIR and modifies types and operations to a84/// form that is appropriate for the specific target. LLVM IR has specific85/// idioms that are used for distinct target processor and ABI combinations.86class TargetRewrite : public fir::impl::TargetRewritePassBase<TargetRewrite> {87public:88 using TargetRewritePassBase<TargetRewrite>::TargetRewritePassBase;89 90 void runOnOperation() override final {91 auto &context = getContext();92 mlir::OpBuilder rewriter(&context);93 94 auto mod = getModule();95 if (!forcedTargetTriple.empty())96 fir::setTargetTriple(mod, forcedTargetTriple);97 98 if (!forcedTargetCPU.empty())99 fir::setTargetCPU(mod, forcedTargetCPU);100 101 if (!forcedTuneCPU.empty())102 fir::setTuneCPU(mod, forcedTuneCPU);103 104 if (!forcedTargetFeatures.empty())105 fir::setTargetFeatures(mod, forcedTargetFeatures);106 107 // TargetRewrite will require querying the type storage sizes, if it was108 // not set already, create a DataLayoutSpec for the ModuleOp now.109 std::optional<mlir::DataLayout> dl =110 fir::support::getOrSetMLIRDataLayout(mod, /*allowDefaultLayout=*/true);111 if (!dl) {112 mlir::emitError(mod.getLoc(),113 "module operation must carry a data layout attribute "114 "to perform target ABI rewrites on FIR");115 signalPassFailure();116 return;117 }118 119 auto specifics = fir::CodeGenSpecifics::get(120 mod.getContext(), fir::getTargetTriple(mod), fir::getKindMapping(mod),121 fir::getTargetCPU(mod), fir::getTargetFeatures(mod), *dl,122 fir::getTuneCPU(mod));123 124 setMembers(specifics.get(), &rewriter, &*dl);125 126 // Perform type conversion on signatures and call sites.127 if (mlir::failed(convertTypes(mod))) {128 mlir::emitError(mlir::UnknownLoc::get(&context),129 "error in converting types to target abi");130 signalPassFailure();131 }132 133 // Convert ops in target-specific patterns.134 mod.walk([&](mlir::Operation *op) {135 if (auto call = mlir::dyn_cast<fir::CallOp>(op)) {136 if (!hasPortableSignature(call.getFunctionType(), op))137 convertCallOp(call, call.getFunctionType());138 } else if (auto dispatch = mlir::dyn_cast<fir::DispatchOp>(op)) {139 if (!hasPortableSignature(dispatch.getFunctionType(), op))140 convertCallOp(dispatch, dispatch.getFunctionType());141 } else if (auto gpuLaunchFunc =142 mlir::dyn_cast<mlir::gpu::LaunchFuncOp>(op)) {143 llvm::SmallVector<mlir::Type> operandsTypes;144 for (auto arg : gpuLaunchFunc.getKernelOperands())145 operandsTypes.push_back(arg.getType());146 auto fctTy = mlir::FunctionType::get(&context, operandsTypes,147 gpuLaunchFunc.getResultTypes());148 if (!hasPortableSignature(fctTy, op))149 convertCallOp(gpuLaunchFunc, fctTy);150 } else if (auto addr = mlir::dyn_cast<fir::AddrOfOp>(op)) {151 if (mlir::isa<mlir::FunctionType>(addr.getType()) &&152 !hasPortableSignature(addr.getType(), op))153 convertAddrOp(addr);154 }155 });156 157 clearMembers();158 }159 160 mlir::ModuleOp getModule() { return getOperation(); }161 162 template <typename Ty, typename Callback>163 std::optional<std::function<mlir::Value(mlir::Operation *)>>164 rewriteCallResultType(mlir::Location loc, mlir::Type originalResTy,165 Ty &newResTys,166 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs,167 Callback &newOpers, mlir::Value &savedStackPtr,168 fir::CodeGenSpecifics::Marshalling &m) {169 // Currently, targets mandate COMPLEX or STRUCT is a single aggregate or170 // packed scalar, including the sret case.171 assert(m.size() == 1 && "return type not supported on this target");172 auto resTy = std::get<mlir::Type>(m[0]);173 auto attr = std::get<fir::CodeGenSpecifics::Attributes>(m[0]);174 if (attr.isSRet()) {175 assert(fir::isa_ref_type(resTy) && "must be a memory reference type");176 // Save the stack pointer, if it has not been saved for this call yet.177 // We will need to restore it after the call, because the alloca178 // needs to be deallocated.179 if (!savedStackPtr)180 savedStackPtr = genStackSave(loc);181 mlir::Value stack =182 fir::AllocaOp::create(*rewriter, loc, fir::dyn_cast_ptrEleTy(resTy));183 newInTyAndAttrs.push_back(m[0]);184 newOpers.push_back(stack);185 return [=](mlir::Operation *) -> mlir::Value {186 auto memTy = fir::ReferenceType::get(originalResTy);187 auto cast = fir::ConvertOp::create(*rewriter, loc, memTy, stack);188 return fir::LoadOp::create(*rewriter, loc, cast);189 };190 }191 newResTys.push_back(resTy);192 return [=, &savedStackPtr](mlir::Operation *call) -> mlir::Value {193 // We are going to generate an alloca, so save the stack pointer.194 if (!savedStackPtr)195 savedStackPtr = genStackSave(loc);196 return this->convertValueInMemory(loc, call->getResult(0), originalResTy,197 /*inputMayBeBigger=*/true);198 };199 }200 201 template <typename Ty, typename Callback>202 std::optional<std::function<mlir::Value(mlir::Operation *)>>203 rewriteCallComplexResultType(204 mlir::Location loc, mlir::ComplexType ty, Ty &newResTys,205 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, Callback &newOpers,206 mlir::Value &savedStackPtr) {207 if (noComplexConversion) {208 newResTys.push_back(ty);209 return std::nullopt;210 }211 auto m = specifics->complexReturnType(loc, ty.getElementType());212 return rewriteCallResultType(loc, ty, newResTys, newInTyAndAttrs, newOpers,213 savedStackPtr, m);214 }215 216 template <typename Ty, typename Callback>217 std::optional<std::function<mlir::Value(mlir::Operation *)>>218 rewriteCallStructResultType(219 mlir::Location loc, fir::RecordType recTy, Ty &newResTys,220 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, Callback &newOpers,221 mlir::Value &savedStackPtr) {222 if (noStructConversion) {223 newResTys.push_back(recTy);224 return std::nullopt;225 }226 auto m = specifics->structReturnType(loc, recTy);227 return rewriteCallResultType(loc, recTy, newResTys, newInTyAndAttrs,228 newOpers, savedStackPtr, m);229 }230 231 void passArgumentOnStackOrWithNewType(232 mlir::Location loc, fir::CodeGenSpecifics::TypeAndAttr newTypeAndAttr,233 mlir::Type oldType, mlir::Value oper,234 llvm::SmallVectorImpl<mlir::Value> &newOpers,235 mlir::Value &savedStackPtr) {236 auto resTy = std::get<mlir::Type>(newTypeAndAttr);237 auto attr = std::get<fir::CodeGenSpecifics::Attributes>(newTypeAndAttr);238 // We are going to generate an alloca, so save the stack pointer.239 if (!savedStackPtr)240 savedStackPtr = genStackSave(loc);241 if (attr.isByVal()) {242 mlir::Value mem = fir::AllocaOp::create(*rewriter, loc, oldType);243 fir::StoreOp::create(*rewriter, loc, oper, mem);244 if (mem.getType() != resTy)245 mem = fir::ConvertOp::create(*rewriter, loc, resTy, mem);246 newOpers.push_back(mem);247 } else {248 mlir::Value bitcast =249 convertValueInMemory(loc, oper, resTy, /*inputMayBeBigger=*/false);250 newOpers.push_back(bitcast);251 }252 }253 254 // Do a bitcast (convert a value via its memory representation).255 // The input and output types may have different storage sizes,256 // "inputMayBeBigger" should be set to indicate which of the input or257 // output type may be bigger in order for the load/store to be safe.258 // The mismatch comes from the fact that the LLVM register used for passing259 // may be bigger than the value being passed (e.g., passing260 // a `!fir.type<t{fir.array<3xi8>}>` into an i32 LLVM register).261 mlir::Value convertValueInMemory(mlir::Location loc, mlir::Value value,262 mlir::Type newType, bool inputMayBeBigger) {263 if (inputMayBeBigger) {264 auto newRefTy = fir::ReferenceType::get(newType);265 auto mem = fir::AllocaOp::create(*rewriter, loc, value.getType());266 fir::StoreOp::create(*rewriter, loc, value, mem);267 auto cast = fir::ConvertOp::create(*rewriter, loc, newRefTy, mem);268 return fir::LoadOp::create(*rewriter, loc, cast);269 } else {270 auto oldRefTy = fir::ReferenceType::get(value.getType());271 auto mem = fir::AllocaOp::create(*rewriter, loc, newType);272 auto cast = fir::ConvertOp::create(*rewriter, loc, oldRefTy, mem);273 fir::StoreOp::create(*rewriter, loc, value, cast);274 return fir::LoadOp::create(*rewriter, loc, mem);275 }276 }277 278 void passSplitArgument(mlir::Location loc,279 fir::CodeGenSpecifics::Marshalling splitArgs,280 mlir::Type oldType, mlir::Value oper,281 llvm::SmallVectorImpl<mlir::Value> &newOpers,282 mlir::Value &savedStackPtr) {283 // COMPLEX or struct argument split into separate arguments284 if (!fir::isa_complex(oldType)) {285 // Cast original operand to a tuple of the new arguments286 // via memory.287 llvm::SmallVector<mlir::Type> partTypes;288 for (auto argPart : splitArgs)289 partTypes.push_back(std::get<mlir::Type>(argPart));290 mlir::Type tupleType =291 mlir::TupleType::get(oldType.getContext(), partTypes);292 if (!savedStackPtr)293 savedStackPtr = genStackSave(loc);294 oper = convertValueInMemory(loc, oper, tupleType,295 /*inputMayBeBigger=*/false);296 }297 auto iTy = rewriter->getIntegerType(32);298 for (auto e : llvm::enumerate(splitArgs)) {299 auto &tup = e.value();300 auto ty = std::get<mlir::Type>(tup);301 auto index = e.index();302 auto idx = rewriter->getIntegerAttr(iTy, index);303 auto val = fir::ExtractValueOp::create(*rewriter, loc, ty, oper,304 rewriter->getArrayAttr(idx));305 newOpers.push_back(val);306 }307 }308 309 void rewriteCallOperands(310 mlir::Location loc, fir::CodeGenSpecifics::Marshalling passArgAs,311 mlir::Type originalArgTy, mlir::Value oper,312 llvm::SmallVectorImpl<mlir::Value> &newOpers, mlir::Value &savedStackPtr,313 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs) {314 if (passArgAs.size() == 1) {315 // COMPLEX or derived type is passed as a single argument.316 passArgumentOnStackOrWithNewType(loc, passArgAs[0], originalArgTy, oper,317 newOpers, savedStackPtr);318 } else {319 // COMPLEX or derived type is split into separate arguments320 passSplitArgument(loc, passArgAs, originalArgTy, oper, newOpers,321 savedStackPtr);322 }323 newInTyAndAttrs.insert(newInTyAndAttrs.end(), passArgAs.begin(),324 passArgAs.end());325 }326 327 template <typename CPLX>328 void rewriteCallComplexInputType(329 mlir::Location loc, CPLX ty, mlir::Value oper,330 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs,331 llvm::SmallVectorImpl<mlir::Value> &newOpers,332 mlir::Value &savedStackPtr) {333 if (noComplexConversion) {334 newInTyAndAttrs.push_back(fir::CodeGenSpecifics::getTypeAndAttr(ty));335 newOpers.push_back(oper);336 return;337 }338 auto m = specifics->complexArgumentType(loc, ty.getElementType());339 rewriteCallOperands(loc, m, ty, oper, newOpers, savedStackPtr,340 newInTyAndAttrs);341 }342 343 void rewriteCallStructInputType(344 mlir::Location loc, fir::RecordType recTy, mlir::Value oper,345 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs,346 llvm::SmallVectorImpl<mlir::Value> &newOpers,347 mlir::Value &savedStackPtr) {348 if (noStructConversion) {349 newInTyAndAttrs.push_back(fir::CodeGenSpecifics::getTypeAndAttr(recTy));350 newOpers.push_back(oper);351 return;352 }353 auto structArgs =354 specifics->structArgumentType(loc, recTy, newInTyAndAttrs);355 rewriteCallOperands(loc, structArgs, recTy, oper, newOpers, savedStackPtr,356 newInTyAndAttrs);357 }358 359 static bool hasByValOrSRetArgs(360 const fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs) {361 return llvm::any_of(newInTyAndAttrs, [](auto arg) {362 const auto &attr = std::get<fir::CodeGenSpecifics::Attributes>(arg);363 return attr.isByVal() || attr.isSRet();364 });365 }366 367 // Convert fir.call and fir.dispatch Ops.368 template <typename A>369 void convertCallOp(A callOp, mlir::FunctionType fnTy) {370 auto loc = callOp.getLoc();371 rewriter->setInsertionPoint(callOp);372 llvm::SmallVector<mlir::Type> newResTys;373 fir::CodeGenSpecifics::Marshalling newInTyAndAttrs;374 llvm::SmallVector<mlir::Value> newOpers;375 mlir::Value savedStackPtr = nullptr;376 377 // If the call is indirect, the first argument must still be the function378 // to call.379 int dropFront = 0;380 if constexpr (std::is_same_v<std::decay_t<A>, fir::CallOp>) {381 if (!callOp.getCallee()) {382 newInTyAndAttrs.push_back(383 fir::CodeGenSpecifics::getTypeAndAttr(fnTy.getInput(0)));384 newOpers.push_back(callOp.getOperand(0));385 dropFront = 1;386 }387 } else if constexpr (std::is_same_v<std::decay_t<A>, fir::DispatchOp>) {388 dropFront = 1; // First operand is the polymorphic object.389 }390 391 // Determine the rewrite function, `wrap`, for the result value.392 std::optional<std::function<mlir::Value(mlir::Operation *)>> wrap;393 if (fnTy.getResults().size() == 1) {394 mlir::Type ty = fnTy.getResult(0);395 llvm::TypeSwitch<mlir::Type>(ty)396 .template Case<mlir::ComplexType>([&](mlir::ComplexType cmplx) {397 wrap = rewriteCallComplexResultType(loc, cmplx, newResTys,398 newInTyAndAttrs, newOpers,399 savedStackPtr);400 })401 .template Case<fir::RecordType>([&](fir::RecordType recTy) {402 wrap = rewriteCallStructResultType(loc, recTy, newResTys,403 newInTyAndAttrs, newOpers,404 savedStackPtr);405 })406 .Default([&](mlir::Type ty) { newResTys.push_back(ty); });407 } else if (fnTy.getResults().size() > 1) {408 TODO(loc, "multiple results not supported yet");409 }410 411 llvm::SmallVector<mlir::Type> trailingInTys;412 llvm::SmallVector<mlir::Value> trailingOpers;413 llvm::SmallVector<mlir::Value> operands;414 unsigned passArgShift = 0;415 if constexpr (std::is_same_v<std::decay_t<A>, mlir::gpu::LaunchFuncOp>)416 operands = callOp.getKernelOperands();417 else418 operands = callOp.getOperands().drop_front(dropFront);419 for (auto e : llvm::enumerate(420 llvm::zip(fnTy.getInputs().drop_front(dropFront), operands))) {421 mlir::Type ty = std::get<0>(e.value());422 mlir::Value oper = std::get<1>(e.value());423 unsigned index = e.index();424 llvm::TypeSwitch<mlir::Type>(ty)425 .template Case<fir::BoxCharType>([&](fir::BoxCharType boxTy) {426 if constexpr (std::is_same_v<std::decay_t<A>, fir::CallOp>) {427 if (noCharacterConversion) {428 newInTyAndAttrs.push_back(429 fir::CodeGenSpecifics::getTypeAndAttr(boxTy));430 newOpers.push_back(oper);431 return;432 }433 } else {434 // TODO: dispatch case; it used to be a to-do because of sret,435 // but is not tested and maybe should be removed. This pass is436 // anyway ran after lowering fir.dispatch in flang, so maybe that437 // should just be a requirement of the pass.438 TODO(loc, "ABI of fir.dispatch with character arguments");439 }440 auto m = specifics->boxcharArgumentType(boxTy.getEleTy());441 auto unbox = fir::UnboxCharOp::create(442 *rewriter, loc, std::get<mlir::Type>(m[0]),443 std::get<mlir::Type>(m[1]), oper);444 // unboxed CHARACTER arguments445 for (auto e : llvm::enumerate(m)) {446 unsigned idx = e.index();447 auto attr =448 std::get<fir::CodeGenSpecifics::Attributes>(e.value());449 auto argTy = std::get<mlir::Type>(e.value());450 if (attr.isAppend()) {451 trailingInTys.push_back(argTy);452 trailingOpers.push_back(unbox.getResult(idx));453 } else {454 newInTyAndAttrs.push_back(e.value());455 newOpers.push_back(unbox.getResult(idx));456 }457 }458 })459 .template Case<mlir::ComplexType>([&](mlir::ComplexType cmplx) {460 rewriteCallComplexInputType(loc, cmplx, oper, newInTyAndAttrs,461 newOpers, savedStackPtr);462 })463 .template Case<fir::RecordType>([&](fir::RecordType recTy) {464 rewriteCallStructInputType(loc, recTy, oper, newInTyAndAttrs,465 newOpers, savedStackPtr);466 })467 .template Case<mlir::TupleType>([&](mlir::TupleType tuple) {468 if (fir::isCharacterProcedureTuple(tuple)) {469 mlir::ModuleOp module = getModule();470 if constexpr (std::is_same_v<std::decay_t<A>, fir::CallOp>) {471 if (callOp.getCallee()) {472 llvm::StringRef charProcAttr =473 fir::getCharacterProcedureDummyAttrName();474 // The charProcAttr attribute is only used as a safety to475 // confirm that this is a dummy procedure and should be split.476 // It cannot be used to match because attributes are not477 // available in case of indirect calls.478 auto funcOp = module.lookupSymbol<mlir::func::FuncOp>(479 *callOp.getCallee());480 if (funcOp &&481 !funcOp.template getArgAttrOfType<mlir::UnitAttr>(482 index, charProcAttr))483 mlir::emitError(loc, "tuple argument will be split even "484 "though it does not have the `" +485 charProcAttr + "` attribute");486 }487 }488 mlir::Type funcPointerType = tuple.getType(0);489 mlir::Type lenType = tuple.getType(1);490 fir::FirOpBuilder builder(*rewriter, module);491 auto [funcPointer, len] =492 fir::factory::extractCharacterProcedureTuple(builder, loc,493 oper);494 newInTyAndAttrs.push_back(495 fir::CodeGenSpecifics::getTypeAndAttr(funcPointerType));496 newOpers.push_back(funcPointer);497 trailingInTys.push_back(lenType);498 trailingOpers.push_back(len);499 } else {500 newInTyAndAttrs.push_back(501 fir::CodeGenSpecifics::getTypeAndAttr(tuple));502 newOpers.push_back(oper);503 }504 })505 .Default([&](mlir::Type ty) {506 if constexpr (std::is_same_v<std::decay_t<A>, fir::DispatchOp>) {507 if (callOp.getPassArgPos() && *callOp.getPassArgPos() == index)508 passArgShift = newOpers.size() - *callOp.getPassArgPos();509 }510 newInTyAndAttrs.push_back(511 fir::CodeGenSpecifics::getTypeAndAttr(ty));512 newOpers.push_back(oper);513 });514 }515 516 llvm::SmallVector<mlir::Type> newInTypes = toTypeList(newInTyAndAttrs);517 newInTypes.insert(newInTypes.end(), trailingInTys.begin(),518 trailingInTys.end());519 newOpers.insert(newOpers.end(), trailingOpers.begin(), trailingOpers.end());520 521 llvm::SmallVector<mlir::Value, 1> newCallResults;522 // TODO propagate/update call argument and result attributes.523 if constexpr (std::is_same_v<std::decay_t<A>, mlir::gpu::LaunchFuncOp>) {524 mlir::Value asyncToken = callOp.getAsyncToken();525 auto newCall = A::create(*rewriter, loc, callOp.getKernel(),526 callOp.getGridSizeOperandValues(),527 callOp.getBlockSizeOperandValues(),528 callOp.getDynamicSharedMemorySize(), newOpers,529 asyncToken ? asyncToken.getType() : nullptr,530 callOp.getAsyncDependencies(),531 /*clusterSize=*/std::nullopt);532 if (callOp.getClusterSizeX())533 newCall.getClusterSizeXMutable().assign(callOp.getClusterSizeX());534 if (callOp.getClusterSizeY())535 newCall.getClusterSizeYMutable().assign(callOp.getClusterSizeY());536 if (callOp.getClusterSizeZ())537 newCall.getClusterSizeZMutable().assign(callOp.getClusterSizeZ());538 newCallResults.append(newCall.result_begin(), newCall.result_end());539 if (auto cudaProcAttr =540 callOp->template getAttrOfType<cuf::ProcAttributeAttr>(541 cuf::getProcAttrName())) {542 newCall->setAttr(cuf::getProcAttrName(), cudaProcAttr);543 }544 } else if constexpr (std::is_same_v<std::decay_t<A>, fir::CallOp>) {545 fir::CallOp newCall;546 if (callOp.getCallee()) {547 newCall = fir::CallOp::create(*rewriter, loc, *callOp.getCallee(),548 newResTys, newOpers);549 } else {550 newOpers[0].setType(mlir::FunctionType::get(551 callOp.getContext(),552 mlir::TypeRange{newInTypes}.drop_front(dropFront), newResTys));553 newCall = fir::CallOp::create(*rewriter, loc, newResTys, newOpers);554 }555 newCall.setFastmathAttr(callOp.getFastmathAttr());556 // Always set ABI argument attributes on call operations, even when557 // direct, as required by558 // https://llvm.org/docs/LangRef.html#parameter-attributes.559 if (hasByValOrSRetArgs(newInTyAndAttrs)) {560 llvm::SmallVector<mlir::Attribute> argAttrsArray;561 for (const auto &arg :562 llvm::ArrayRef<fir::CodeGenSpecifics::TypeAndAttr>(newInTyAndAttrs)563 .drop_front(dropFront)) {564 mlir::NamedAttrList argAttrs;565 const auto &attr = std::get<fir::CodeGenSpecifics::Attributes>(arg);566 if (attr.isByVal()) {567 mlir::Type elemType =568 fir::dyn_cast_ptrOrBoxEleTy(std::get<mlir::Type>(arg));569 argAttrs.set(mlir::LLVM::LLVMDialect::getByValAttrName(),570 mlir::TypeAttr::get(elemType));571 } else if (attr.isSRet()) {572 mlir::Type elemType =573 fir::dyn_cast_ptrOrBoxEleTy(std::get<mlir::Type>(arg));574 argAttrs.set(mlir::LLVM::LLVMDialect::getStructRetAttrName(),575 mlir::TypeAttr::get(elemType));576 }577 if (auto align = attr.getAlignment()) {578 argAttrs.set(579 mlir::LLVM::LLVMDialect::getAlignAttrName(),580 rewriter->getIntegerAttr(rewriter->getIntegerType(32), align));581 }582 argAttrsArray.emplace_back(583 argAttrs.getDictionary(rewriter->getContext()));584 }585 newCall.setArgAttrsAttr(rewriter->getArrayAttr(argAttrsArray));586 }587 LLVM_DEBUG(llvm::dbgs() << "replacing call with " << newCall << '\n');588 if (wrap)589 newCallResults.push_back((*wrap)(newCall.getOperation()));590 else591 newCallResults.append(newCall.result_begin(), newCall.result_end());592 } else {593 fir::DispatchOp dispatchOp = A::create(594 *rewriter, loc, newResTys,595 rewriter->getStringAttr(callOp.getMethod()), callOp.getOperands()[0],596 newOpers,597 rewriter->getI32IntegerAttr(*callOp.getPassArgPos() + passArgShift),598 /*arg_attrs=*/nullptr, /*res_attrs=*/nullptr,599 callOp.getProcedureAttrsAttr());600 if (wrap)601 newCallResults.push_back((*wrap)(dispatchOp.getOperation()));602 else603 newCallResults.append(dispatchOp.result_begin(),604 dispatchOp.result_end());605 }606 607 if (newCallResults.size() <= 1) {608 if (savedStackPtr) {609 if (newCallResults.size() == 1) {610 // We assume that all the allocas are inserted before611 // the operation that defines the new call result.612 rewriter->setInsertionPointAfterValue(newCallResults[0]);613 } else {614 // If the call does not have results, then insert615 // stack restore after the original call operation.616 rewriter->setInsertionPointAfter(callOp);617 }618 genStackRestore(loc, savedStackPtr);619 }620 replaceOp(callOp, newCallResults);621 } else {622 // The TODO is duplicated here to make sure this part623 // handles the stackrestore insertion properly, if624 // we add support for multiple call results.625 TODO(loc, "multiple results not supported yet");626 }627 }628 629 // Result type fixup for ComplexType.630 template <typename Ty>631 void lowerComplexSignatureRes(632 mlir::Location loc, mlir::ComplexType cmplx, Ty &newResTys,633 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs) {634 if (noComplexConversion) {635 newResTys.push_back(cmplx);636 return;637 }638 for (auto &tup :639 specifics->complexReturnType(loc, cmplx.getElementType())) {640 auto argTy = std::get<mlir::Type>(tup);641 if (std::get<fir::CodeGenSpecifics::Attributes>(tup).isSRet())642 newInTyAndAttrs.push_back(tup);643 else644 newResTys.push_back(argTy);645 }646 }647 648 // Argument type fixup for ComplexType.649 void lowerComplexSignatureArg(650 mlir::Location loc, mlir::ComplexType cmplx,651 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs) {652 if (noComplexConversion) {653 newInTyAndAttrs.push_back(fir::CodeGenSpecifics::getTypeAndAttr(cmplx));654 } else {655 auto cplxArgs =656 specifics->complexArgumentType(loc, cmplx.getElementType());657 newInTyAndAttrs.insert(newInTyAndAttrs.end(), cplxArgs.begin(),658 cplxArgs.end());659 }660 }661 662 template <typename Ty>663 void664 lowerStructSignatureRes(mlir::Location loc, fir::RecordType recTy,665 Ty &newResTys,666 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs) {667 if (noComplexConversion) {668 newResTys.push_back(recTy);669 return;670 } else {671 for (auto &tup : specifics->structReturnType(loc, recTy)) {672 if (std::get<fir::CodeGenSpecifics::Attributes>(tup).isSRet())673 newInTyAndAttrs.push_back(tup);674 else675 newResTys.push_back(std::get<mlir::Type>(tup));676 }677 }678 }679 680 void681 lowerStructSignatureArg(mlir::Location loc, fir::RecordType recTy,682 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs) {683 if (noStructConversion) {684 newInTyAndAttrs.push_back(fir::CodeGenSpecifics::getTypeAndAttr(recTy));685 return;686 }687 auto structArgs =688 specifics->structArgumentType(loc, recTy, newInTyAndAttrs);689 newInTyAndAttrs.insert(newInTyAndAttrs.end(), structArgs.begin(),690 structArgs.end());691 }692 693 llvm::SmallVector<mlir::Type>694 toTypeList(const fir::CodeGenSpecifics::Marshalling &marshalled) {695 llvm::SmallVector<mlir::Type> typeList;696 for (auto &typeAndAttr : marshalled)697 typeList.emplace_back(std::get<mlir::Type>(typeAndAttr));698 return typeList;699 }700 701 /// Taking the address of a function. Modify the signature as needed.702 void convertAddrOp(fir::AddrOfOp addrOp) {703 rewriter->setInsertionPoint(addrOp);704 auto addrTy = mlir::cast<mlir::FunctionType>(addrOp.getType());705 fir::CodeGenSpecifics::Marshalling newInTyAndAttrs;706 llvm::SmallVector<mlir::Type> newResTys;707 auto loc = addrOp.getLoc();708 for (mlir::Type ty : addrTy.getResults()) {709 llvm::TypeSwitch<mlir::Type>(ty)710 .Case<mlir::ComplexType>([&](mlir::ComplexType ty) {711 lowerComplexSignatureRes(loc, ty, newResTys, newInTyAndAttrs);712 })713 .Case<fir::RecordType>([&](fir::RecordType ty) {714 lowerStructSignatureRes(loc, ty, newResTys, newInTyAndAttrs);715 })716 .Default([&](mlir::Type ty) { newResTys.push_back(ty); });717 }718 llvm::SmallVector<mlir::Type> trailingInTys;719 for (mlir::Type ty : addrTy.getInputs()) {720 llvm::TypeSwitch<mlir::Type>(ty)721 .Case<fir::BoxCharType>([&](auto box) {722 if (noCharacterConversion) {723 newInTyAndAttrs.push_back(724 fir::CodeGenSpecifics::getTypeAndAttr(box));725 } else {726 for (auto &tup : specifics->boxcharArgumentType(box.getEleTy())) {727 auto attr = std::get<fir::CodeGenSpecifics::Attributes>(tup);728 auto argTy = std::get<mlir::Type>(tup);729 if (attr.isAppend())730 trailingInTys.push_back(argTy);731 else732 newInTyAndAttrs.push_back(tup);733 }734 }735 })736 .Case<mlir::ComplexType>([&](mlir::ComplexType ty) {737 lowerComplexSignatureArg(loc, ty, newInTyAndAttrs);738 })739 .Case<mlir::TupleType>([&](mlir::TupleType tuple) {740 if (fir::isCharacterProcedureTuple(tuple)) {741 newInTyAndAttrs.push_back(742 fir::CodeGenSpecifics::getTypeAndAttr(tuple.getType(0)));743 trailingInTys.push_back(tuple.getType(1));744 } else {745 newInTyAndAttrs.push_back(746 fir::CodeGenSpecifics::getTypeAndAttr(ty));747 }748 })749 .template Case<fir::RecordType>([&](fir::RecordType recTy) {750 lowerStructSignatureArg(loc, recTy, newInTyAndAttrs);751 })752 .Default([&](mlir::Type ty) {753 newInTyAndAttrs.push_back(754 fir::CodeGenSpecifics::getTypeAndAttr(ty));755 });756 }757 llvm::SmallVector<mlir::Type> newInTypes = toTypeList(newInTyAndAttrs);758 // append trailing input types759 newInTypes.insert(newInTypes.end(), trailingInTys.begin(),760 trailingInTys.end());761 // replace this op with a new one with the updated signature762 auto newTy = rewriter->getFunctionType(newInTypes, newResTys);763 auto newOp = fir::AddrOfOp::create(*rewriter, addrOp.getLoc(), newTy,764 addrOp.getSymbol());765 replaceOp(addrOp, newOp.getResult());766 }767 768 /// Convert the type signatures on all the functions present in the module.769 /// As the type signature is being changed, this must also update the770 /// function itself to use any new arguments, etc.771 llvm::LogicalResult convertTypes(mlir::ModuleOp mod) {772 mlir::MLIRContext *ctx = mod->getContext();773 auto targetCPU = specifics->getTargetCPU();774 mlir::StringAttr targetCPUAttr =775 targetCPU.empty() ? nullptr : mlir::StringAttr::get(ctx, targetCPU);776 auto tuneCPU = specifics->getTuneCPU();777 mlir::StringAttr tuneCPUAttr =778 tuneCPU.empty() ? nullptr : mlir::StringAttr::get(ctx, tuneCPU);779 auto targetFeaturesAttr = specifics->getTargetFeatures();780 781 for (auto fn : mod.getOps<mlir::func::FuncOp>()) {782 if (targetCPUAttr)783 fn->setAttr("target_cpu", targetCPUAttr);784 785 if (tuneCPUAttr)786 fn->setAttr("tune_cpu", tuneCPUAttr);787 788 if (targetFeaturesAttr)789 fn->setAttr("target_features", targetFeaturesAttr);790 791 convertSignature<mlir::func::ReturnOp, mlir::func::FuncOp>(fn);792 }793 794 for (auto gpuMod : mod.getOps<mlir::gpu::GPUModuleOp>()) {795 for (auto fn : gpuMod.getOps<mlir::func::FuncOp>())796 convertSignature<mlir::func::ReturnOp, mlir::func::FuncOp>(fn);797 for (auto fn : gpuMod.getOps<mlir::gpu::GPUFuncOp>())798 convertSignature<mlir::gpu::ReturnOp, mlir::gpu::GPUFuncOp>(fn);799 }800 801 return mlir::success();802 }803 804 // Returns true if the function should be interoperable with C.805 static bool isFuncWithCCallingConvention(mlir::Operation *op) {806 auto funcOp = mlir::dyn_cast<mlir::func::FuncOp>(op);807 if (!funcOp)808 return false;809 return op->hasAttrOfType<mlir::UnitAttr>(810 fir::FIROpsDialect::getFirRuntimeAttrName()) ||811 op->hasAttrOfType<mlir::StringAttr>(fir::getSymbolAttrName());812 }813 814 /// If the signature does not need any special target-specific conversions,815 /// then it is considered portable for any target, and this function will816 /// return `true`. Otherwise, the signature is not portable and `false` is817 /// returned.818 bool hasPortableSignature(mlir::Type signature, mlir::Operation *op) {819 assert(mlir::isa<mlir::FunctionType>(signature));820 auto func = mlir::dyn_cast<mlir::FunctionType>(signature);821 bool hasCCallingConv = isFuncWithCCallingConvention(op);822 for (auto ty : func.getResults())823 if ((mlir::isa<fir::BoxCharType>(ty) && !noCharacterConversion) ||824 (fir::isa_complex(ty) && !noComplexConversion) ||825 (mlir::isa<mlir::IntegerType>(ty) && hasCCallingConv) ||826 (mlir::isa<fir::RecordType>(ty) && !noStructConversion)) {827 LLVM_DEBUG(llvm::dbgs() << "rewrite " << signature << " for target\n");828 return false;829 }830 for (auto ty : func.getInputs())831 if (((mlir::isa<fir::BoxCharType>(ty) ||832 fir::isCharacterProcedureTuple(ty)) &&833 !noCharacterConversion) ||834 (fir::isa_complex(ty) && !noComplexConversion) ||835 (mlir::isa<mlir::IntegerType>(ty) && hasCCallingConv) ||836 (mlir::isa<fir::RecordType>(ty) && !noStructConversion)) {837 LLVM_DEBUG(llvm::dbgs() << "rewrite " << signature << " for target\n");838 return false;839 }840 return true;841 }842 843 /// Determine if the signature has host associations. The host association844 /// argument may need special target specific rewriting.845 template <typename OpTy>846 static bool hasHostAssociations(OpTy func) {847 std::size_t end = func.getFunctionType().getInputs().size();848 for (std::size_t i = 0; i < end; ++i)849 if (func.template getArgAttrOfType<mlir::UnitAttr>(850 i, fir::getHostAssocAttrName()))851 return true;852 return false;853 }854 855 /// Rewrite the signatures and body of the `FuncOp`s in the module for856 /// the immediately subsequent target code gen.857 template <typename ReturnOpTy, typename FuncOpTy>858 void convertSignature(FuncOpTy func) {859 auto funcTy = mlir::cast<mlir::FunctionType>(func.getFunctionType());860 if (hasPortableSignature(funcTy, func) && !hasHostAssociations(func))861 return;862 llvm::SmallVector<mlir::Type> newResTys;863 fir::CodeGenSpecifics::Marshalling newInTyAndAttrs;864 llvm::SmallVector<std::pair<unsigned, mlir::NamedAttribute>> savedAttrs;865 llvm::SmallVector<std::pair<unsigned, mlir::NamedAttribute>> extraAttrs;866 llvm::SmallVector<FixupTy> fixups;867 llvm::SmallVector<std::pair<unsigned, mlir::NamedAttrList>, 1> resultAttrs;868 869 // Save argument attributes in case there is a shift so we can replace them870 // correctly.871 for (auto e : llvm::enumerate(funcTy.getInputs())) {872 unsigned index = e.index();873 llvm::ArrayRef<mlir::NamedAttribute> attrs =874 mlir::function_interface_impl::getArgAttrs(func, index);875 for (mlir::NamedAttribute attr : attrs) {876 savedAttrs.push_back({index, attr});877 }878 }879 880 // Count the number of arguments that have to stay in place at the end of881 // the argument list.882 unsigned trailingArgs = 0;883 if constexpr (std::is_same_v<FuncOpTy, mlir::gpu::GPUFuncOp>) {884 trailingArgs =885 func.getNumWorkgroupAttributions() + func.getNumPrivateAttributions();886 }887 888 // Convert return value(s)889 for (auto ty : funcTy.getResults())890 llvm::TypeSwitch<mlir::Type>(ty)891 .template Case<mlir::ComplexType>([&](mlir::ComplexType cmplx) {892 if (noComplexConversion)893 newResTys.push_back(cmplx);894 else895 doComplexReturn(func, cmplx, newResTys, newInTyAndAttrs, fixups);896 })897 .template Case<mlir::IntegerType>([&](mlir::IntegerType intTy) {898 auto m = specifics->integerArgumentType(func.getLoc(), intTy);899 assert(m.size() == 1);900 auto attr = std::get<fir::CodeGenSpecifics::Attributes>(m[0]);901 auto retTy = std::get<mlir::Type>(m[0]);902 std::size_t resId = newResTys.size();903 llvm::StringRef extensionAttrName = attr.getIntExtensionAttrName();904 if (!extensionAttrName.empty() &&905 isFuncWithCCallingConvention(func))906 resultAttrs.emplace_back(907 resId, rewriter->getNamedAttr(extensionAttrName,908 rewriter->getUnitAttr()));909 newResTys.push_back(retTy);910 })911 .template Case<fir::RecordType>([&](fir::RecordType recTy) {912 doStructReturn(func, recTy, newResTys, newInTyAndAttrs, fixups);913 })914 .Default([&](mlir::Type ty) { newResTys.push_back(ty); });915 916 // Saved potential shift in argument. Handling of result can add arguments917 // at the beginning of the function signature.918 unsigned argumentShift = newInTyAndAttrs.size();919 920 // Convert arguments921 llvm::SmallVector<mlir::Type> trailingTys;922 for (auto e : llvm::enumerate(funcTy.getInputs())) {923 auto ty = e.value();924 unsigned index = e.index();925 llvm::TypeSwitch<mlir::Type>(ty)926 .template Case<fir::BoxCharType>([&](fir::BoxCharType boxTy) {927 if (noCharacterConversion) {928 newInTyAndAttrs.push_back(929 fir::CodeGenSpecifics::getTypeAndAttr(boxTy));930 } else {931 // Convert a CHARACTER argument type. This can involve separating932 // the pointer and the LEN into two arguments and moving the LEN933 // argument to the end of the arg list.934 for (auto &tup :935 specifics->boxcharArgumentType(boxTy.getEleTy())) {936 auto attr = std::get<fir::CodeGenSpecifics::Attributes>(tup);937 auto argTy = std::get<mlir::Type>(tup);938 if (attr.isAppend()) {939 trailingTys.push_back(argTy);940 } else {941 fixups.emplace_back(FixupTy::Codes::Trailing,942 newInTyAndAttrs.size(),943 trailingTys.size());944 newInTyAndAttrs.push_back(tup);945 }946 }947 }948 })949 .template Case<mlir::ComplexType>([&](mlir::ComplexType cmplx) {950 doComplexArg(func, cmplx, newInTyAndAttrs, fixups);951 })952 .template Case<mlir::TupleType>([&](mlir::TupleType tuple) {953 if (fir::isCharacterProcedureTuple(tuple)) {954 fixups.emplace_back(FixupTy::Codes::TrailingCharProc,955 newInTyAndAttrs.size(), trailingTys.size());956 newInTyAndAttrs.push_back(957 fir::CodeGenSpecifics::getTypeAndAttr(tuple.getType(0)));958 trailingTys.push_back(tuple.getType(1));959 } else {960 newInTyAndAttrs.push_back(961 fir::CodeGenSpecifics::getTypeAndAttr(ty));962 }963 })964 .template Case<mlir::IntegerType>([&](mlir::IntegerType intTy) {965 auto m = specifics->integerArgumentType(func.getLoc(), intTy);966 assert(m.size() == 1);967 auto attr = std::get<fir::CodeGenSpecifics::Attributes>(m[0]);968 auto argNo = newInTyAndAttrs.size();969 llvm::StringRef extensionAttrName = attr.getIntExtensionAttrName();970 if (!extensionAttrName.empty() &&971 isFuncWithCCallingConvention(func))972 fixups.emplace_back(FixupTy::Codes::ArgumentType, argNo,973 [=](FuncOpTy func) {974 func.setArgAttr(975 argNo, extensionAttrName,976 mlir::UnitAttr::get(func.getContext()));977 });978 979 newInTyAndAttrs.push_back(m[0]);980 })981 .template Case<fir::RecordType>([&](fir::RecordType recTy) {982 doStructArg(func, recTy, newInTyAndAttrs, fixups);983 })984 .Default([&](mlir::Type ty) {985 newInTyAndAttrs.push_back(986 fir::CodeGenSpecifics::getTypeAndAttr(ty));987 });988 989 if (func.template getArgAttrOfType<mlir::UnitAttr>(990 index, fir::getHostAssocAttrName())) {991 extraAttrs.push_back(992 {newInTyAndAttrs.size() - 1,993 rewriter->getNamedAttr("llvm.nest", rewriter->getUnitAttr())});994 }995 }996 997 // Add the argument at the end if the number of trailing arguments is 0,998 // otherwise insert the argument at the appropriate index.999 auto addOrInsertArgument = [&](mlir::Type ty, mlir::Location loc) {1000 unsigned inputIndex = func.front().getArguments().size() - trailingArgs;1001 auto newArg = trailingArgs == 01002 ? func.front().addArgument(ty, loc)1003 : func.front().insertArgument(inputIndex, ty, loc);1004 return newArg;1005 };1006 1007 if (!func.empty()) {1008 // If the function has a body, then apply the fixups to the arguments and1009 // return ops as required. These fixups are done in place.1010 auto loc = func.getLoc();1011 const auto fixupSize = fixups.size();1012 const auto oldArgTys = func.getFunctionType().getInputs();1013 int offset = 0;1014 for (std::remove_const_t<decltype(fixupSize)> i = 0; i < fixupSize; ++i) {1015 const auto &fixup = fixups[i];1016 mlir::Type fixupType =1017 fixup.index < newInTyAndAttrs.size()1018 ? std::get<mlir::Type>(newInTyAndAttrs[fixup.index])1019 : mlir::Type{};1020 switch (fixup.code) {1021 case FixupTy::Codes::ArgumentAsLoad: {1022 // Argument was pass-by-value, but is now pass-by-reference and1023 // possibly with a different element type.1024 auto newArg =1025 func.front().insertArgument(fixup.index, fixupType, loc);1026 rewriter->setInsertionPointToStart(&func.front());1027 auto oldArgTy =1028 fir::ReferenceType::get(oldArgTys[fixup.index - offset]);1029 auto cast = fir::ConvertOp::create(*rewriter, loc, oldArgTy, newArg);1030 auto load = fir::LoadOp::create(*rewriter, loc, cast);1031 func.getArgument(fixup.index + 1).replaceAllUsesWith(load);1032 func.front().eraseArgument(fixup.index + 1);1033 } break;1034 case FixupTy::Codes::ArgumentType: {1035 // Argument is pass-by-value, but its type has likely been modified to1036 // suit the target ABI convention.1037 auto oldArgTy = oldArgTys[fixup.index - offset];1038 // If type did not change, keep the original argument.1039 if (fixupType == oldArgTy)1040 break;1041 1042 auto newArg =1043 func.front().insertArgument(fixup.index, fixupType, loc);1044 rewriter->setInsertionPointToStart(&func.front());1045 mlir::Value bitcast = convertValueInMemory(loc, newArg, oldArgTy,1046 /*inputMayBeBigger=*/true);1047 func.getArgument(fixup.index + 1).replaceAllUsesWith(bitcast);1048 func.front().eraseArgument(fixup.index + 1);1049 LLVM_DEBUG(llvm::dbgs()1050 << "old argument: " << oldArgTy << ", repl: " << bitcast1051 << ", new argument: "1052 << func.getArgument(fixup.index).getType() << '\n');1053 } break;1054 case FixupTy::Codes::CharPair: {1055 // The FIR boxchar argument has been split into a pair of distinct1056 // arguments that are in juxtaposition to each other.1057 auto newArg =1058 func.front().insertArgument(fixup.index, fixupType, loc);1059 if (fixup.second == 1) {1060 rewriter->setInsertionPointToStart(&func.front());1061 auto boxTy = oldArgTys[fixup.index - offset - fixup.second];1062 auto box = fir::EmboxCharOp::create(1063 *rewriter, loc, boxTy,1064 func.front().getArgument(fixup.index - 1), newArg);1065 func.getArgument(fixup.index + 1).replaceAllUsesWith(box);1066 func.front().eraseArgument(fixup.index + 1);1067 offset++;1068 }1069 } break;1070 case FixupTy::Codes::ReturnAsStore: {1071 // The value being returned is now being returned in memory (callee1072 // stack space) through a hidden reference argument.1073 auto newArg =1074 func.front().insertArgument(fixup.index, fixupType, loc);1075 offset++;1076 func.walk([&](ReturnOpTy ret) {1077 rewriter->setInsertionPoint(ret);1078 auto oldOper = ret.getOperand(0);1079 auto oldOperTy = fir::ReferenceType::get(oldOper.getType());1080 auto cast =1081 fir::ConvertOp::create(*rewriter, loc, oldOperTy, newArg);1082 fir::StoreOp::create(*rewriter, loc, oldOper, cast);1083 ReturnOpTy::create(*rewriter, loc);1084 ret.erase();1085 });1086 } break;1087 case FixupTy::Codes::ReturnType: {1088 // The function is still returning a value, but its type has likely1089 // changed to suit the target ABI convention.1090 func.walk([&](ReturnOpTy ret) {1091 rewriter->setInsertionPoint(ret);1092 auto oldOper = ret.getOperand(0);1093 mlir::Value bitcast =1094 convertValueInMemory(loc, oldOper, newResTys[fixup.index],1095 /*inputMayBeBigger=*/false);1096 ReturnOpTy::create(*rewriter, loc, bitcast);1097 ret.erase();1098 });1099 } break;1100 case FixupTy::Codes::Split: {1101 // The FIR argument has been split into a pair of distinct arguments1102 // that are in juxtaposition to each other. (For COMPLEX value or1103 // derived type passed with VALUE in BIND(C) context).1104 auto newArg =1105 func.front().insertArgument(fixup.index, fixupType, loc);1106 if (fixup.second == 1) {1107 rewriter->setInsertionPointToStart(&func.front());1108 mlir::Value firstArg = func.front().getArgument(fixup.index - 1);1109 mlir::Type originalTy =1110 oldArgTys[fixup.index - offset - fixup.second];1111 mlir::Type pairTy = originalTy;1112 if (!fir::isa_complex(originalTy)) {1113 pairTy = mlir::TupleType::get(1114 originalTy.getContext(),1115 mlir::TypeRange{firstArg.getType(), newArg.getType()});1116 }1117 auto undef = fir::UndefOp::create(*rewriter, loc, pairTy);1118 auto iTy = rewriter->getIntegerType(32);1119 auto zero = rewriter->getIntegerAttr(iTy, 0);1120 auto one = rewriter->getIntegerAttr(iTy, 1);1121 mlir::Value pair1 = fir::InsertValueOp::create(1122 *rewriter, loc, pairTy, undef, firstArg,1123 rewriter->getArrayAttr(zero));1124 mlir::Value pair =1125 fir::InsertValueOp::create(*rewriter, loc, pairTy, pair1,1126 newArg, rewriter->getArrayAttr(one));1127 // Cast local argument tuple to original type via memory if needed.1128 if (pairTy != originalTy)1129 pair = convertValueInMemory(loc, pair, originalTy,1130 /*inputMayBeBigger=*/true);1131 func.getArgument(fixup.index + 1).replaceAllUsesWith(pair);1132 func.front().eraseArgument(fixup.index + 1);1133 offset++;1134 }1135 } break;1136 case FixupTy::Codes::Trailing: {1137 // The FIR argument has been split into a pair of distinct arguments.1138 // The first part of the pair appears in the original argument1139 // position. The second part of the pair is appended after all the1140 // original arguments. (Boxchar arguments.)1141 auto newBufArg =1142 func.front().insertArgument(fixup.index, fixupType, loc);1143 auto newLenArg = addOrInsertArgument(trailingTys[fixup.second], loc);1144 auto boxTy = oldArgTys[fixup.index - offset];1145 rewriter->setInsertionPointToStart(&func.front());1146 auto box = fir::EmboxCharOp::create(*rewriter, loc, boxTy, newBufArg,1147 newLenArg);1148 func.getArgument(fixup.index + 1).replaceAllUsesWith(box);1149 func.front().eraseArgument(fixup.index + 1);1150 } break;1151 case FixupTy::Codes::TrailingCharProc: {1152 // The FIR character procedure argument tuple must be split into a1153 // pair of distinct arguments. The first part of the pair appears in1154 // the original argument position. The second part of the pair is1155 // appended after all the original arguments.1156 auto newProcPointerArg =1157 func.front().insertArgument(fixup.index, fixupType, loc);1158 auto newLenArg = addOrInsertArgument(trailingTys[fixup.second], loc);1159 auto tupleType = oldArgTys[fixup.index - offset];1160 rewriter->setInsertionPointToStart(&func.front());1161 fir::FirOpBuilder builder(*rewriter, getModule());1162 auto tuple = fir::factory::createCharacterProcedureTuple(1163 builder, loc, tupleType, newProcPointerArg, newLenArg);1164 func.getArgument(fixup.index + 1).replaceAllUsesWith(tuple);1165 func.front().eraseArgument(fixup.index + 1);1166 } break;1167 }1168 }1169 }1170 1171 llvm::SmallVector<mlir::Type> newInTypes = toTypeList(newInTyAndAttrs);1172 // Set the new type and finalize the arguments, etc.1173 newInTypes.insert(newInTypes.end(), trailingTys.begin(), trailingTys.end());1174 auto newFuncTy =1175 mlir::FunctionType::get(func.getContext(), newInTypes, newResTys);1176 LLVM_DEBUG(llvm::dbgs() << "new func: " << newFuncTy << '\n');1177 func.setType(newFuncTy);1178 1179 for (std::pair<unsigned, mlir::NamedAttribute> extraAttr : extraAttrs)1180 func.setArgAttr(extraAttr.first, extraAttr.second.getName(),1181 extraAttr.second.getValue());1182 1183 for (auto [resId, resAttrList] : resultAttrs)1184 for (mlir::NamedAttribute resAttr : resAttrList)1185 func.setResultAttr(resId, resAttr.getName(), resAttr.getValue());1186 1187 // Replace attributes to the correct argument if there was an argument shift1188 // to the right.1189 if (argumentShift > 0) {1190 for (std::pair<unsigned, mlir::NamedAttribute> savedAttr : savedAttrs) {1191 func.removeArgAttr(savedAttr.first, savedAttr.second.getName());1192 func.setArgAttr(savedAttr.first + argumentShift,1193 savedAttr.second.getName(),1194 savedAttr.second.getValue());1195 }1196 }1197 1198 for (auto &fixup : fixups) {1199 if constexpr (std::is_same_v<FuncOpTy, mlir::func::FuncOp>)1200 if (fixup.finalizer)1201 (*fixup.finalizer)(func);1202 if constexpr (std::is_same_v<FuncOpTy, mlir::gpu::GPUFuncOp>)1203 if (fixup.gpuFinalizer)1204 (*fixup.gpuFinalizer)(func);1205 }1206 }1207 1208 template <typename OpTy, typename Ty, typename FIXUPS>1209 void doReturn(OpTy func, Ty &newResTys,1210 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs,1211 FIXUPS &fixups, fir::CodeGenSpecifics::Marshalling &m) {1212 assert(m.size() == 1 &&1213 "expect result to be turned into single argument or result so far");1214 auto &tup = m[0];1215 auto attr = std::get<fir::CodeGenSpecifics::Attributes>(tup);1216 auto argTy = std::get<mlir::Type>(tup);1217 if (attr.isSRet()) {1218 unsigned argNo = newInTyAndAttrs.size();1219 if (auto align = attr.getAlignment())1220 fixups.emplace_back(1221 FixupTy::Codes::ReturnAsStore, argNo, [=](OpTy func) {1222 auto elemType = fir::dyn_cast_ptrOrBoxEleTy(1223 func.getFunctionType().getInput(argNo));1224 func.setArgAttr(argNo, "llvm.sret",1225 mlir::TypeAttr::get(elemType));1226 func.setArgAttr(argNo, "llvm.align",1227 rewriter->getIntegerAttr(1228 rewriter->getIntegerType(32), align));1229 });1230 else1231 fixups.emplace_back(FixupTy::Codes::ReturnAsStore, argNo,1232 [=](OpTy func) {1233 auto elemType = fir::dyn_cast_ptrOrBoxEleTy(1234 func.getFunctionType().getInput(argNo));1235 func.setArgAttr(argNo, "llvm.sret",1236 mlir::TypeAttr::get(elemType));1237 });1238 newInTyAndAttrs.push_back(tup);1239 return;1240 }1241 if (auto align = attr.getAlignment())1242 fixups.emplace_back(1243 FixupTy::Codes::ReturnType, newResTys.size(), [=](OpTy func) {1244 func.setArgAttr(1245 newResTys.size(), "llvm.align",1246 rewriter->getIntegerAttr(rewriter->getIntegerType(32), align));1247 });1248 else1249 fixups.emplace_back(FixupTy::Codes::ReturnType, newResTys.size());1250 newResTys.push_back(argTy);1251 }1252 1253 /// Convert a complex return value. This can involve converting the return1254 /// value to a "hidden" first argument or packing the complex into a wide1255 /// GPR.1256 template <typename OpTy, typename Ty, typename FIXUPS>1257 void doComplexReturn(OpTy func, mlir::ComplexType cmplx, Ty &newResTys,1258 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs,1259 FIXUPS &fixups) {1260 if (noComplexConversion) {1261 newResTys.push_back(cmplx);1262 return;1263 }1264 auto m =1265 specifics->complexReturnType(func.getLoc(), cmplx.getElementType());1266 doReturn(func, newResTys, newInTyAndAttrs, fixups, m);1267 }1268 1269 template <typename OpTy, typename Ty, typename FIXUPS>1270 void doStructReturn(OpTy func, fir::RecordType recTy, Ty &newResTys,1271 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs,1272 FIXUPS &fixups) {1273 if (noStructConversion) {1274 newResTys.push_back(recTy);1275 return;1276 }1277 auto m = specifics->structReturnType(func.getLoc(), recTy);1278 doReturn(func, newResTys, newInTyAndAttrs, fixups, m);1279 }1280 1281 template <typename OpTy, typename FIXUPS>1282 void createFuncOpArgFixups(1283 OpTy func, fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs,1284 fir::CodeGenSpecifics::Marshalling &argsInTys, FIXUPS &fixups) {1285 const auto fixupCode = argsInTys.size() > 1 ? FixupTy::Codes::Split1286 : FixupTy::Codes::ArgumentType;1287 for (auto e : llvm::enumerate(argsInTys)) {1288 auto &tup = e.value();1289 auto index = e.index();1290 auto attr = std::get<fir::CodeGenSpecifics::Attributes>(tup);1291 auto argNo = newInTyAndAttrs.size();1292 if (attr.isByVal()) {1293 if (auto align = attr.getAlignment())1294 fixups.emplace_back(FixupTy::Codes::ArgumentAsLoad, argNo,1295 [=](OpTy func) {1296 auto elemType = fir::dyn_cast_ptrOrBoxEleTy(1297 func.getFunctionType().getInput(argNo));1298 func.setArgAttr(argNo, "llvm.byval",1299 mlir::TypeAttr::get(elemType));1300 func.setArgAttr(1301 argNo, "llvm.align",1302 rewriter->getIntegerAttr(1303 rewriter->getIntegerType(32), align));1304 });1305 else1306 fixups.emplace_back(FixupTy::Codes::ArgumentAsLoad,1307 newInTyAndAttrs.size(), [=](OpTy func) {1308 auto elemType = fir::dyn_cast_ptrOrBoxEleTy(1309 func.getFunctionType().getInput(argNo));1310 func.setArgAttr(argNo, "llvm.byval",1311 mlir::TypeAttr::get(elemType));1312 });1313 } else {1314 if (auto align = attr.getAlignment())1315 fixups.emplace_back(1316 fixupCode, argNo, index, [=](OpTy func) {1317 func.setArgAttr(argNo, "llvm.align",1318 rewriter->getIntegerAttr(1319 rewriter->getIntegerType(32), align));1320 });1321 else1322 fixups.emplace_back(fixupCode, argNo, index);1323 }1324 newInTyAndAttrs.push_back(tup);1325 }1326 }1327 1328 /// Convert a complex argument value. This can involve storing the value to1329 /// a temporary memory location or factoring the value into two distinct1330 /// arguments.1331 template <typename OpTy, typename FIXUPS>1332 void doComplexArg(OpTy func, mlir::ComplexType cmplx,1333 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs,1334 FIXUPS &fixups) {1335 if (noComplexConversion) {1336 newInTyAndAttrs.push_back(fir::CodeGenSpecifics::getTypeAndAttr(cmplx));1337 return;1338 }1339 auto cplxArgs =1340 specifics->complexArgumentType(func.getLoc(), cmplx.getElementType());1341 createFuncOpArgFixups(func, newInTyAndAttrs, cplxArgs, fixups);1342 }1343 1344 template <typename OpTy, typename FIXUPS>1345 void doStructArg(OpTy func, fir::RecordType recTy,1346 fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs,1347 FIXUPS &fixups) {1348 if (noStructConversion) {1349 newInTyAndAttrs.push_back(fir::CodeGenSpecifics::getTypeAndAttr(recTy));1350 return;1351 }1352 auto structArgs =1353 specifics->structArgumentType(func.getLoc(), recTy, newInTyAndAttrs);1354 createFuncOpArgFixups(func, newInTyAndAttrs, structArgs, fixups);1355 }1356 1357private:1358 // Replace `op` and remove it.1359 void replaceOp(mlir::Operation *op, mlir::ValueRange newValues) {1360 llvm::SmallVector<mlir::Value> casts;1361 for (auto [oldValue, newValue] : llvm::zip(op->getResults(), newValues)) {1362 if (oldValue.getType() == newValue.getType())1363 casts.push_back(newValue);1364 else1365 casts.push_back(fir::ConvertOp::create(*rewriter, op->getLoc(),1366 oldValue.getType(), newValue));1367 }1368 op->replaceAllUsesWith(casts);1369 op->dropAllReferences();1370 op->erase();1371 }1372 1373 inline void setMembers(fir::CodeGenSpecifics *s, mlir::OpBuilder *r,1374 mlir::DataLayout *dl) {1375 specifics = s;1376 rewriter = r;1377 dataLayout = dl;1378 }1379 1380 inline void clearMembers() { setMembers(nullptr, nullptr, nullptr); }1381 1382 // Inserts a call to llvm.stacksave at the current insertion1383 // point and the given location. Returns the call's result Value.1384 inline mlir::Value genStackSave(mlir::Location loc) {1385 fir::FirOpBuilder builder(*rewriter, getModule());1386 return builder.genStackSave(loc);1387 }1388 1389 // Inserts a call to llvm.stackrestore at the current insertion1390 // point and the given location and argument.1391 inline void genStackRestore(mlir::Location loc, mlir::Value sp) {1392 fir::FirOpBuilder builder(*rewriter, getModule());1393 return builder.genStackRestore(loc, sp);1394 }1395 1396 fir::CodeGenSpecifics *specifics = nullptr;1397 mlir::OpBuilder *rewriter = nullptr;1398 mlir::DataLayout *dataLayout = nullptr;1399};1400} // namespace1401