| //===-- TargetRewrite.cpp -------------------------------------------------===// | |
| // | |
| // 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 | |
| // | |
| //===----------------------------------------------------------------------===// | |
| // | |
| // Target rewrite: rewriting of ops to make target-specific lowerings manifest. | |
| // LLVM expects different lowering idioms to be used for distinct target | |
| // triples. These distinctions are handled by this pass. | |
| // | |
| // Coding style: https://mlir.llvm.org/getting_started/DeveloperGuide/ | |
| // | |
| //===----------------------------------------------------------------------===// | |
| #include "flang/Optimizer/CodeGen/CodeGen.h" | |
| #include "flang/Optimizer/Builder/Character.h" | |
| #include "flang/Optimizer/Builder/FIRBuilder.h" | |
| #include "flang/Optimizer/Builder/Todo.h" | |
| #include "flang/Optimizer/CodeGen/Target.h" | |
| #include "flang/Optimizer/Dialect/FIRDialect.h" | |
| #include "flang/Optimizer/Dialect/FIROps.h" | |
| #include "flang/Optimizer/Dialect/FIROpsSupport.h" | |
| #include "flang/Optimizer/Dialect/FIRType.h" | |
| #include "flang/Optimizer/Dialect/Support/FIRContext.h" | |
| #include "flang/Optimizer/Support/DataLayout.h" | |
| #include "mlir/Dialect/DLTI/DLTI.h" | |
| #include "mlir/Dialect/GPU/IR/GPUDialect.h" | |
| #include "mlir/Dialect/LLVMIR/LLVMDialect.h" | |
| #include "mlir/Transforms/DialectConversion.h" | |
| #include "llvm/ADT/STLExtras.h" | |
| #include "llvm/ADT/TypeSwitch.h" | |
| #include "llvm/Support/Debug.h" | |
| #include <optional> | |
| namespace fir { | |
| #define GEN_PASS_DEF_TARGETREWRITEPASS | |
| #include "flang/Optimizer/CodeGen/CGPasses.h.inc" | |
| } // namespace fir | |
| #define DEBUG_TYPE "flang-target-rewrite" | |
| namespace { | |
| /// Fixups for updating a FuncOp's arguments and return values. | |
| struct FixupTy { | |
| enum class Codes { | |
| ArgumentAsLoad, | |
| ArgumentType, | |
| CharPair, | |
| ReturnAsStore, | |
| ReturnType, | |
| Split, | |
| Trailing, | |
| TrailingCharProc | |
| }; | |
| FixupTy(Codes code, std::size_t index, std::size_t second = 0) | |
| : code{code}, index{index}, second{second} {} | |
| FixupTy(Codes code, std::size_t index, | |
| std::function<void(mlir::func::FuncOp)> &&finalizer) | |
| : code{code}, index{index}, finalizer{finalizer} {} | |
| FixupTy(Codes code, std::size_t index, | |
| std::function<void(mlir::gpu::GPUFuncOp)> &&finalizer) | |
| : code{code}, index{index}, gpuFinalizer{finalizer} {} | |
| FixupTy(Codes code, std::size_t index, std::size_t second, | |
| std::function<void(mlir::func::FuncOp)> &&finalizer) | |
| : code{code}, index{index}, second{second}, finalizer{finalizer} {} | |
| FixupTy(Codes code, std::size_t index, std::size_t second, | |
| std::function<void(mlir::gpu::GPUFuncOp)> &&finalizer) | |
| : code{code}, index{index}, second{second}, gpuFinalizer{finalizer} {} | |
| Codes code; | |
| std::size_t index; | |
| std::size_t second{}; | |
| std::optional<std::function<void(mlir::func::FuncOp)>> finalizer{}; | |
| std::optional<std::function<void(mlir::gpu::GPUFuncOp)>> gpuFinalizer{}; | |
| }; // namespace | |
| /// Target-specific rewriting of the FIR. This is a prerequisite pass to code | |
| /// generation that traverses the FIR and modifies types and operations to a | |
| /// form that is appropriate for the specific target. LLVM IR has specific | |
| /// idioms that are used for distinct target processor and ABI combinations. | |
| class TargetRewrite : public fir::impl::TargetRewritePassBase<TargetRewrite> { | |
| public: | |
| using TargetRewritePassBase<TargetRewrite>::TargetRewritePassBase; | |
| void runOnOperation() override final { | |
| auto &context = getContext(); | |
| mlir::OpBuilder rewriter(&context); | |
| auto mod = getModule(); | |
| if (!forcedTargetTriple.empty()) | |
| fir::setTargetTriple(mod, forcedTargetTriple); | |
| if (!forcedTargetCPU.empty()) | |
| fir::setTargetCPU(mod, forcedTargetCPU); | |
| if (!forcedTuneCPU.empty()) | |
| fir::setTuneCPU(mod, forcedTuneCPU); | |
| if (!forcedTargetFeatures.empty()) | |
| fir::setTargetFeatures(mod, forcedTargetFeatures); | |
| if (!forcedTargetABI.empty()) | |
| fir::setTargetABI(mod, forcedTargetABI); | |
| // If no data layout is set, derive it from the target triple to get | |
| // correct target-specific alignments rather than a generic default. | |
| if (!mod.getDataLayoutSpec()) { | |
| llvm::Triple triple(fir::getTargetTriple(mod)); | |
| std::string dlStr = triple.computeDataLayout(); | |
| if (!dlStr.empty()) { | |
| llvm::DataLayout llvmDL(dlStr); | |
| fir::support::setMLIRDataLayout(mod, llvmDL); | |
| } | |
| } | |
| // TargetRewrite will require querying the type storage sizes, if it was | |
| // not set already, create a DataLayoutSpec for the ModuleOp now. | |
| std::optional<mlir::DataLayout> dl = | |
| fir::support::getOrSetMLIRDataLayout(mod, /*allowDefaultLayout=*/true); | |
| if (!dl) { | |
| mlir::emitError(mod.getLoc(), | |
| "module operation must carry a data layout attribute " | |
| "to perform target ABI rewrites on FIR"); | |
| signalPassFailure(); | |
| return; | |
| } | |
| auto specifics = fir::CodeGenSpecifics::get( | |
| mod.getContext(), fir::getTargetTriple(mod), fir::getKindMapping(mod), | |
| fir::getTargetCPU(mod), fir::getTargetFeatures(mod), | |
| fir::getTargetABI(mod), *dl, fir::getTuneCPU(mod)); | |
| setMembers(specifics.get(), &rewriter, &*dl); | |
| // Perform type conversion on signatures and call sites. | |
| if (mlir::failed(convertTypes(mod))) { | |
| mlir::emitError(mlir::UnknownLoc::get(&context), | |
| "error in converting types to target abi"); | |
| signalPassFailure(); | |
| } | |
| // Convert ops in target-specific patterns. | |
| mod.walk([&](mlir::Operation *op) { | |
| if (auto call = mlir::dyn_cast<fir::CallOp>(op)) { | |
| if (!hasPortableSignature(call.getFunctionType(), op)) | |
| convertCallOp(call, call.getFunctionType()); | |
| } else if (auto dispatch = mlir::dyn_cast<fir::DispatchOp>(op)) { | |
| if (!hasPortableSignature(dispatch.getFunctionType(), op)) | |
| convertCallOp(dispatch, dispatch.getFunctionType()); | |
| } else if (auto gpuLaunchFunc = | |
| mlir::dyn_cast<mlir::gpu::LaunchFuncOp>(op)) { | |
| llvm::SmallVector<mlir::Type> operandsTypes; | |
| for (auto arg : gpuLaunchFunc.getKernelOperands()) | |
| operandsTypes.push_back(arg.getType()); | |
| auto fctTy = mlir::FunctionType::get(&context, operandsTypes, | |
| gpuLaunchFunc.getResultTypes()); | |
| if (!hasPortableSignature(fctTy, op)) | |
| convertCallOp(gpuLaunchFunc, fctTy); | |
| } else if (auto addr = mlir::dyn_cast<fir::AddrOfOp>(op)) { | |
| if (mlir::isa<mlir::FunctionType>(addr.getType()) && | |
| !hasPortableSignature(addr.getType(), op)) | |
| convertAddrOp(addr); | |
| } | |
| }); | |
| clearMembers(); | |
| } | |
| mlir::ModuleOp getModule() { return getOperation(); } | |
| template <typename Ty, typename Callback> | |
| std::optional<std::function<mlir::Value(mlir::Operation *)>> | |
| rewriteCallResultType(mlir::Location loc, mlir::Type originalResTy, | |
| Ty &newResTys, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, | |
| Callback &newOpers, mlir::Value &savedStackPtr, | |
| fir::CodeGenSpecifics::Marshalling &m) { | |
| // Currently, targets mandate COMPLEX or STRUCT is a single aggregate or | |
| // packed scalar, including the sret case. | |
| assert(m.size() == 1 && "return type not supported on this target"); | |
| auto resTy = std::get<mlir::Type>(m[0]); | |
| auto attr = std::get<fir::CodeGenSpecifics::Attributes>(m[0]); | |
| if (attr.isSRet()) { | |
| assert(fir::isa_ref_type(resTy) && "must be a memory reference type"); | |
| // Save the stack pointer, if it has not been saved for this call yet. | |
| // We will need to restore it after the call, because the alloca | |
| // needs to be deallocated. | |
| if (!savedStackPtr) | |
| savedStackPtr = genStackSave(loc); | |
| mlir::Value stack = | |
| fir::AllocaOp::create(*rewriter, loc, fir::dyn_cast_ptrEleTy(resTy)); | |
| newInTyAndAttrs.push_back(m[0]); | |
| newOpers.push_back(stack); | |
| return [=](mlir::Operation *) -> mlir::Value { | |
| auto memTy = fir::ReferenceType::get(originalResTy); | |
| auto cast = fir::ConvertOp::create(*rewriter, loc, memTy, stack); | |
| return fir::LoadOp::create(*rewriter, loc, cast); | |
| }; | |
| } | |
| newResTys.push_back(resTy); | |
| return [=, &savedStackPtr](mlir::Operation *call) -> mlir::Value { | |
| // We are going to generate an alloca, so save the stack pointer. | |
| if (!savedStackPtr) | |
| savedStackPtr = genStackSave(loc); | |
| return this->convertValueInMemory(loc, call->getResult(0), originalResTy, | |
| /*inputMayBeBigger=*/true); | |
| }; | |
| } | |
| template <typename Ty, typename Callback> | |
| std::optional<std::function<mlir::Value(mlir::Operation *)>> | |
| rewriteCallComplexResultType( | |
| mlir::Location loc, mlir::ComplexType ty, Ty &newResTys, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, Callback &newOpers, | |
| mlir::Value &savedStackPtr) { | |
| if (noComplexConversion) { | |
| newResTys.push_back(ty); | |
| return std::nullopt; | |
| } | |
| auto m = specifics->complexReturnType(loc, ty.getElementType()); | |
| return rewriteCallResultType(loc, ty, newResTys, newInTyAndAttrs, newOpers, | |
| savedStackPtr, m); | |
| } | |
| template <typename Ty, typename Callback> | |
| std::optional<std::function<mlir::Value(mlir::Operation *)>> | |
| rewriteCallStructResultType( | |
| mlir::Location loc, fir::RecordType recTy, Ty &newResTys, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, Callback &newOpers, | |
| mlir::Value &savedStackPtr) { | |
| if (noStructConversion) { | |
| newResTys.push_back(recTy); | |
| return std::nullopt; | |
| } | |
| auto m = specifics->structReturnType(loc, recTy); | |
| return rewriteCallResultType(loc, recTy, newResTys, newInTyAndAttrs, | |
| newOpers, savedStackPtr, m); | |
| } | |
| void passArgumentOnStackOrWithNewType( | |
| mlir::Location loc, fir::CodeGenSpecifics::TypeAndAttr newTypeAndAttr, | |
| mlir::Type oldType, mlir::Value oper, | |
| llvm::SmallVectorImpl<mlir::Value> &newOpers, | |
| mlir::Value &savedStackPtr) { | |
| auto resTy = std::get<mlir::Type>(newTypeAndAttr); | |
| auto attr = std::get<fir::CodeGenSpecifics::Attributes>(newTypeAndAttr); | |
| // We are going to generate an alloca, so save the stack pointer. | |
| if (!savedStackPtr) | |
| savedStackPtr = genStackSave(loc); | |
| if (attr.isByVal() || attr.isIndirect()) { | |
| // Both forms require the caller to materialize a copy of the value and | |
| // pass the address of that copy. They differ only in the attribute that | |
| // is attached to the argument in the signature. | |
| mlir::Value mem = fir::AllocaOp::create(*rewriter, loc, oldType); | |
| fir::StoreOp::create(*rewriter, loc, oper, mem); | |
| if (mem.getType() != resTy) | |
| mem = fir::ConvertOp::create(*rewriter, loc, resTy, mem); | |
| newOpers.push_back(mem); | |
| } else { | |
| mlir::Value bitcast = | |
| convertValueInMemory(loc, oper, resTy, /*inputMayBeBigger=*/false); | |
| newOpers.push_back(bitcast); | |
| } | |
| } | |
| // Do a bitcast (convert a value via its memory representation). | |
| // The input and output types may have different storage sizes, | |
| // "inputMayBeBigger" should be set to indicate which of the input or | |
| // output type may be bigger in order for the load/store to be safe. | |
| // The mismatch comes from the fact that the LLVM register used for passing | |
| // may be bigger than the value being passed (e.g., passing | |
| // a `!fir.type<t{fir.array<3xi8>}>` into an i32 LLVM register). | |
| mlir::Value convertValueInMemory(mlir::Location loc, mlir::Value value, | |
| mlir::Type newType, bool inputMayBeBigger) { | |
| if (inputMayBeBigger) { | |
| auto newRefTy = fir::ReferenceType::get(newType); | |
| auto mem = fir::AllocaOp::create(*rewriter, loc, value.getType()); | |
| fir::StoreOp::create(*rewriter, loc, value, mem); | |
| auto cast = fir::ConvertOp::create(*rewriter, loc, newRefTy, mem); | |
| return fir::LoadOp::create(*rewriter, loc, cast); | |
| } else { | |
| auto oldRefTy = fir::ReferenceType::get(value.getType()); | |
| auto mem = fir::AllocaOp::create(*rewriter, loc, newType); | |
| auto cast = fir::ConvertOp::create(*rewriter, loc, oldRefTy, mem); | |
| fir::StoreOp::create(*rewriter, loc, value, cast); | |
| return fir::LoadOp::create(*rewriter, loc, mem); | |
| } | |
| } | |
| void passSplitArgument(mlir::Location loc, | |
| fir::CodeGenSpecifics::Marshalling splitArgs, | |
| mlir::Type oldType, mlir::Value oper, | |
| llvm::SmallVectorImpl<mlir::Value> &newOpers, | |
| mlir::Value &savedStackPtr) { | |
| // COMPLEX or struct argument split into separate arguments | |
| if (!fir::isa_complex(oldType)) { | |
| // Cast original operand to a tuple of the new arguments | |
| // via memory. | |
| llvm::SmallVector<mlir::Type> partTypes; | |
| for (auto argPart : splitArgs) | |
| partTypes.push_back(std::get<mlir::Type>(argPart)); | |
| mlir::Type tupleType = | |
| mlir::TupleType::get(oldType.getContext(), partTypes); | |
| if (!savedStackPtr) | |
| savedStackPtr = genStackSave(loc); | |
| oper = convertValueInMemory(loc, oper, tupleType, | |
| /*inputMayBeBigger=*/false); | |
| } | |
| auto iTy = rewriter->getIntegerType(32); | |
| for (auto e : llvm::enumerate(splitArgs)) { | |
| auto &tup = e.value(); | |
| auto ty = std::get<mlir::Type>(tup); | |
| auto index = e.index(); | |
| auto idx = rewriter->getIntegerAttr(iTy, index); | |
| auto val = fir::ExtractValueOp::create(*rewriter, loc, ty, oper, | |
| rewriter->getArrayAttr(idx)); | |
| newOpers.push_back(val); | |
| } | |
| } | |
| void rewriteCallOperands( | |
| mlir::Location loc, fir::CodeGenSpecifics::Marshalling passArgAs, | |
| mlir::Type originalArgTy, mlir::Value oper, | |
| llvm::SmallVectorImpl<mlir::Value> &newOpers, mlir::Value &savedStackPtr, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs) { | |
| if (passArgAs.size() == 1) { | |
| // COMPLEX or derived type is passed as a single argument. | |
| passArgumentOnStackOrWithNewType(loc, passArgAs[0], originalArgTy, oper, | |
| newOpers, savedStackPtr); | |
| } else { | |
| // COMPLEX or derived type is split into separate arguments | |
| passSplitArgument(loc, passArgAs, originalArgTy, oper, newOpers, | |
| savedStackPtr); | |
| } | |
| newInTyAndAttrs.insert(newInTyAndAttrs.end(), passArgAs.begin(), | |
| passArgAs.end()); | |
| } | |
| template <typename CPLX> | |
| void rewriteCallComplexInputType( | |
| mlir::Location loc, CPLX ty, mlir::Value oper, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, | |
| llvm::SmallVectorImpl<mlir::Value> &newOpers, | |
| mlir::Value &savedStackPtr) { | |
| if (noComplexConversion) { | |
| newInTyAndAttrs.push_back(fir::CodeGenSpecifics::getTypeAndAttr(ty)); | |
| newOpers.push_back(oper); | |
| return; | |
| } | |
| auto m = specifics->complexArgumentType(loc, ty.getElementType()); | |
| rewriteCallOperands(loc, m, ty, oper, newOpers, savedStackPtr, | |
| newInTyAndAttrs); | |
| } | |
| void rewriteCallStructInputType( | |
| mlir::Location loc, fir::RecordType recTy, mlir::Value oper, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, | |
| llvm::SmallVectorImpl<mlir::Value> &newOpers, | |
| mlir::Value &savedStackPtr) { | |
| if (noStructConversion) { | |
| newInTyAndAttrs.push_back(fir::CodeGenSpecifics::getTypeAndAttr(recTy)); | |
| newOpers.push_back(oper); | |
| return; | |
| } | |
| auto structArgs = | |
| specifics->structArgumentType(loc, recTy, newInTyAndAttrs); | |
| rewriteCallOperands(loc, structArgs, recTy, oper, newOpers, savedStackPtr, | |
| newInTyAndAttrs); | |
| } | |
| static bool hasAbiArgAttributes( | |
| const fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs) { | |
| return llvm::any_of(newInTyAndAttrs, [](auto arg) { | |
| const auto &attr = std::get<fir::CodeGenSpecifics::Attributes>(arg); | |
| return attr.isByVal() || attr.isSRet() || attr.isIndirect(); | |
| }); | |
| } | |
| // Convert fir.call and fir.dispatch Ops. | |
| template <typename A> | |
| void convertCallOp(A callOp, mlir::FunctionType fnTy) { | |
| auto loc = callOp.getLoc(); | |
| rewriter->setInsertionPoint(callOp); | |
| llvm::SmallVector<mlir::Type> newResTys; | |
| fir::CodeGenSpecifics::Marshalling newInTyAndAttrs; | |
| llvm::SmallVector<mlir::Value> newOpers; | |
| mlir::Value savedStackPtr = nullptr; | |
| // If the call is indirect, the first argument must still be the function | |
| // to call. | |
| int dropFront = 0; | |
| if constexpr (std::is_same_v<std::decay_t<A>, fir::CallOp>) { | |
| if (!callOp.getCallee()) { | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(fnTy.getInput(0))); | |
| newOpers.push_back(callOp.getOperand(0)); | |
| dropFront = 1; | |
| } | |
| } else if constexpr (std::is_same_v<std::decay_t<A>, fir::DispatchOp>) { | |
| dropFront = 1; // First operand is the polymorphic object. | |
| } | |
| // Determine the rewrite function, `wrap`, for the result value. | |
| std::optional<std::function<mlir::Value(mlir::Operation *)>> wrap; | |
| if (fnTy.getResults().size() == 1) { | |
| mlir::Type ty = fnTy.getResult(0); | |
| llvm::TypeSwitch<mlir::Type>(ty) | |
| .Case([&](mlir::ComplexType cmplx) { | |
| wrap = rewriteCallComplexResultType(loc, cmplx, newResTys, | |
| newInTyAndAttrs, newOpers, | |
| savedStackPtr); | |
| }) | |
| .Case([&](fir::RecordType recTy) { | |
| wrap = rewriteCallStructResultType(loc, recTy, newResTys, | |
| newInTyAndAttrs, newOpers, | |
| savedStackPtr); | |
| }) | |
| .Default([&](mlir::Type ty) { newResTys.push_back(ty); }); | |
| } else if (fnTy.getResults().size() > 1) { | |
| TODO(loc, "multiple results not supported yet"); | |
| } | |
| llvm::SmallVector<mlir::Type> trailingInTys; | |
| llvm::SmallVector<mlir::Value> trailingOpers; | |
| llvm::SmallVector<mlir::Value> operands; | |
| unsigned passArgShift = 0; | |
| if constexpr (std::is_same_v<std::decay_t<A>, mlir::gpu::LaunchFuncOp>) | |
| operands = callOp.getKernelOperands(); | |
| else | |
| operands = callOp.getOperands().drop_front(dropFront); | |
| for (auto e : llvm::enumerate( | |
| llvm::zip(fnTy.getInputs().drop_front(dropFront), operands))) { | |
| mlir::Type ty = std::get<0>(e.value()); | |
| mlir::Value oper = std::get<1>(e.value()); | |
| unsigned index = e.index(); | |
| llvm::TypeSwitch<mlir::Type>(ty) | |
| .Case([&](fir::BoxCharType boxTy) { | |
| if constexpr (std::is_same_v<std::decay_t<A>, fir::CallOp>) { | |
| if (noCharacterConversion) { | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(boxTy)); | |
| newOpers.push_back(oper); | |
| return; | |
| } | |
| } else { | |
| // TODO: dispatch case; it used to be a to-do because of sret, | |
| // but is not tested and maybe should be removed. This pass is | |
| // anyway ran after lowering fir.dispatch in flang, so maybe that | |
| // should just be a requirement of the pass. | |
| TODO(loc, "ABI of fir.dispatch with character arguments"); | |
| } | |
| auto m = specifics->boxcharArgumentType(boxTy.getEleTy()); | |
| auto unbox = fir::UnboxCharOp::create( | |
| *rewriter, loc, std::get<mlir::Type>(m[0]), | |
| std::get<mlir::Type>(m[1]), oper); | |
| // unboxed CHARACTER arguments | |
| for (auto e : llvm::enumerate(m)) { | |
| unsigned idx = e.index(); | |
| auto attr = | |
| std::get<fir::CodeGenSpecifics::Attributes>(e.value()); | |
| auto argTy = std::get<mlir::Type>(e.value()); | |
| if (attr.isAppend()) { | |
| trailingInTys.push_back(argTy); | |
| trailingOpers.push_back(unbox.getResult(idx)); | |
| } else { | |
| newInTyAndAttrs.push_back(e.value()); | |
| newOpers.push_back(unbox.getResult(idx)); | |
| } | |
| } | |
| }) | |
| .Case([&](mlir::ComplexType cmplx) { | |
| rewriteCallComplexInputType(loc, cmplx, oper, newInTyAndAttrs, | |
| newOpers, savedStackPtr); | |
| }) | |
| .Case([&](fir::RecordType recTy) { | |
| rewriteCallStructInputType(loc, recTy, oper, newInTyAndAttrs, | |
| newOpers, savedStackPtr); | |
| }) | |
| .Case([&](mlir::TupleType tuple) { | |
| if (fir::isCharacterProcedureTuple(tuple)) { | |
| mlir::ModuleOp module = getModule(); | |
| if constexpr (std::is_same_v<std::decay_t<A>, fir::CallOp>) { | |
| if (callOp.getCallee()) { | |
| llvm::StringRef charProcAttr = | |
| fir::getCharacterProcedureDummyAttrName(); | |
| // The charProcAttr attribute is only used as a safety to | |
| // confirm that this is a dummy procedure and should be split. | |
| // It cannot be used to match because attributes are not | |
| // available in case of indirect calls. | |
| auto funcOp = module.lookupSymbol<mlir::func::FuncOp>( | |
| *callOp.getCallee()); | |
| if (funcOp && | |
| !funcOp.template getArgAttrOfType<mlir::UnitAttr>( | |
| index, charProcAttr)) | |
| mlir::emitError(loc, "tuple argument will be split even " | |
| "though it does not have the `" + | |
| charProcAttr + "` attribute"); | |
| } | |
| } | |
| mlir::Type funcPointerType = tuple.getType(0); | |
| mlir::Type lenType = tuple.getType(1); | |
| fir::FirOpBuilder builder(*rewriter, module); | |
| auto [funcPointer, len] = | |
| fir::factory::extractCharacterProcedureTuple(builder, loc, | |
| oper); | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(funcPointerType)); | |
| newOpers.push_back(funcPointer); | |
| trailingInTys.push_back(lenType); | |
| trailingOpers.push_back(len); | |
| } else { | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(tuple)); | |
| newOpers.push_back(oper); | |
| } | |
| }) | |
| .Default([&](mlir::Type ty) { | |
| if constexpr (std::is_same_v<std::decay_t<A>, fir::DispatchOp>) { | |
| if (callOp.getPassArgPos() && *callOp.getPassArgPos() == index) | |
| passArgShift = newOpers.size() - *callOp.getPassArgPos(); | |
| } | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(ty)); | |
| newOpers.push_back(oper); | |
| }); | |
| } | |
| llvm::SmallVector<mlir::Type> newInTypes = toTypeList(newInTyAndAttrs); | |
| newInTypes.insert(newInTypes.end(), trailingInTys.begin(), | |
| trailingInTys.end()); | |
| newOpers.insert(newOpers.end(), trailingOpers.begin(), trailingOpers.end()); | |
| llvm::SmallVector<mlir::Value, 1> newCallResults; | |
| // TODO propagate/update call argument and result attributes. | |
| if constexpr (std::is_same_v<std::decay_t<A>, mlir::gpu::LaunchFuncOp>) { | |
| mlir::Value asyncToken = callOp.getAsyncToken(); | |
| auto newCall = A::create( | |
| *rewriter, loc, callOp.getKernel(), callOp.getGridSizeOperandValues(), | |
| callOp.getBlockSizeOperandValues(), | |
| callOp.getDynamicSharedMemorySize(), newOpers, | |
| asyncToken ? asyncToken.getType() : nullptr, | |
| callOp.getAsyncDependencies(), callOp.getAsyncObject(), | |
| /*clusterSize=*/std::nullopt); | |
| if (callOp.getClusterSizeX()) | |
| newCall.getClusterSizeXMutable().assign(callOp.getClusterSizeX()); | |
| if (callOp.getClusterSizeY()) | |
| newCall.getClusterSizeYMutable().assign(callOp.getClusterSizeY()); | |
| if (callOp.getClusterSizeZ()) | |
| newCall.getClusterSizeZMutable().assign(callOp.getClusterSizeZ()); | |
| newCallResults.append(newCall.result_begin(), newCall.result_end()); | |
| if (auto cudaProcAttr = | |
| callOp->template getAttrOfType<cuf::ProcAttributeAttr>( | |
| cuf::getProcAttrName())) { | |
| newCall->setAttr(cuf::getProcAttrName(), cudaProcAttr); | |
| } | |
| } else if constexpr (std::is_same_v<std::decay_t<A>, fir::CallOp>) { | |
| fir::CallOp newCall; | |
| if (callOp.getCallee()) { | |
| newCall = fir::CallOp::create(*rewriter, loc, *callOp.getCallee(), | |
| newResTys, newOpers); | |
| } else { | |
| newOpers[0].setType(mlir::FunctionType::get( | |
| callOp.getContext(), | |
| mlir::TypeRange{newInTypes}.drop_front(dropFront), newResTys)); | |
| newCall = fir::CallOp::create(*rewriter, loc, newResTys, newOpers, | |
| fir::CallOp::Properties{}); | |
| } | |
| newCall.setFastmathAttr(callOp.getFastmathAttr()); | |
| // Always set ABI argument attributes on call operations, even when | |
| // direct, as required by | |
| // https://llvm.org/docs/LangRef.html#parameter-attributes. | |
| if (hasAbiArgAttributes(newInTyAndAttrs)) { | |
| llvm::SmallVector<mlir::Attribute> argAttrsArray; | |
| for (const auto &arg : | |
| llvm::ArrayRef<fir::CodeGenSpecifics::TypeAndAttr>(newInTyAndAttrs) | |
| .drop_front(dropFront)) { | |
| mlir::NamedAttrList argAttrs; | |
| const auto &attr = std::get<fir::CodeGenSpecifics::Attributes>(arg); | |
| if (attr.isByVal()) { | |
| mlir::Type elemType = | |
| fir::dyn_cast_ptrOrBoxEleTy(std::get<mlir::Type>(arg)); | |
| argAttrs.set(mlir::LLVM::LLVMDialect::getByValAttrName(), | |
| mlir::TypeAttr::get(elemType)); | |
| } else if (attr.isSRet()) { | |
| mlir::Type elemType = | |
| fir::dyn_cast_ptrOrBoxEleTy(std::get<mlir::Type>(arg)); | |
| argAttrs.set(mlir::LLVM::LLVMDialect::getStructRetAttrName(), | |
| mlir::TypeAttr::get(elemType)); | |
| } | |
| if (auto align = attr.getAlignment()) { | |
| argAttrs.set( | |
| mlir::LLVM::LLVMDialect::getAlignAttrName(), | |
| rewriter->getIntegerAttr(rewriter->getIntegerType(32), align)); | |
| } | |
| argAttrsArray.emplace_back( | |
| argAttrs.getDictionary(rewriter->getContext())); | |
| } | |
| newCall.setArgAttrsAttr(rewriter->getArrayAttr(argAttrsArray)); | |
| } | |
| LLVM_DEBUG(llvm::dbgs() << "replacing call with " << newCall << '\n'); | |
| if (wrap) | |
| newCallResults.push_back((*wrap)(newCall.getOperation())); | |
| else | |
| newCallResults.append(newCall.result_begin(), newCall.result_end()); | |
| } else { | |
| fir::DispatchOp dispatchOp = A::create( | |
| *rewriter, loc, newResTys, | |
| rewriter->getStringAttr(callOp.getMethod()), callOp.getOperands()[0], | |
| newOpers, | |
| rewriter->getI32IntegerAttr(*callOp.getPassArgPos() + passArgShift), | |
| /*arg_attrs=*/nullptr, /*res_attrs=*/nullptr, | |
| callOp.getProcedureAttrsAttr()); | |
| if (wrap) | |
| newCallResults.push_back((*wrap)(dispatchOp.getOperation())); | |
| else | |
| newCallResults.append(dispatchOp.result_begin(), | |
| dispatchOp.result_end()); | |
| } | |
| if (newCallResults.size() <= 1) { | |
| if (savedStackPtr) { | |
| if (newCallResults.size() == 1) { | |
| // We assume that all the allocas are inserted before | |
| // the operation that defines the new call result. | |
| rewriter->setInsertionPointAfterValue(newCallResults[0]); | |
| } else { | |
| // If the call does not have results, then insert | |
| // stack restore after the original call operation. | |
| rewriter->setInsertionPointAfter(callOp); | |
| } | |
| genStackRestore(loc, savedStackPtr); | |
| } | |
| replaceOp(callOp, newCallResults); | |
| } else { | |
| // The TODO is duplicated here to make sure this part | |
| // handles the stackrestore insertion properly, if | |
| // we add support for multiple call results. | |
| TODO(loc, "multiple results not supported yet"); | |
| } | |
| } | |
| // Result type fixup for ComplexType. | |
| template <typename Ty> | |
| void lowerComplexSignatureRes( | |
| mlir::Location loc, mlir::ComplexType cmplx, Ty &newResTys, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs) { | |
| if (noComplexConversion) { | |
| newResTys.push_back(cmplx); | |
| return; | |
| } | |
| for (auto &tup : | |
| specifics->complexReturnType(loc, cmplx.getElementType())) { | |
| auto argTy = std::get<mlir::Type>(tup); | |
| if (std::get<fir::CodeGenSpecifics::Attributes>(tup).isSRet()) | |
| newInTyAndAttrs.push_back(tup); | |
| else | |
| newResTys.push_back(argTy); | |
| } | |
| } | |
| // Argument type fixup for ComplexType. | |
| void lowerComplexSignatureArg( | |
| mlir::Location loc, mlir::ComplexType cmplx, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs) { | |
| if (noComplexConversion) { | |
| newInTyAndAttrs.push_back(fir::CodeGenSpecifics::getTypeAndAttr(cmplx)); | |
| } else { | |
| auto cplxArgs = | |
| specifics->complexArgumentType(loc, cmplx.getElementType()); | |
| newInTyAndAttrs.insert(newInTyAndAttrs.end(), cplxArgs.begin(), | |
| cplxArgs.end()); | |
| } | |
| } | |
| template <typename Ty> | |
| void | |
| lowerStructSignatureRes(mlir::Location loc, fir::RecordType recTy, | |
| Ty &newResTys, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs) { | |
| if (noComplexConversion) { | |
| newResTys.push_back(recTy); | |
| return; | |
| } else { | |
| for (auto &tup : specifics->structReturnType(loc, recTy)) { | |
| if (std::get<fir::CodeGenSpecifics::Attributes>(tup).isSRet()) | |
| newInTyAndAttrs.push_back(tup); | |
| else | |
| newResTys.push_back(std::get<mlir::Type>(tup)); | |
| } | |
| } | |
| } | |
| void | |
| lowerStructSignatureArg(mlir::Location loc, fir::RecordType recTy, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs) { | |
| if (noStructConversion) { | |
| newInTyAndAttrs.push_back(fir::CodeGenSpecifics::getTypeAndAttr(recTy)); | |
| return; | |
| } | |
| auto structArgs = | |
| specifics->structArgumentType(loc, recTy, newInTyAndAttrs); | |
| newInTyAndAttrs.insert(newInTyAndAttrs.end(), structArgs.begin(), | |
| structArgs.end()); | |
| } | |
| llvm::SmallVector<mlir::Type> | |
| toTypeList(const fir::CodeGenSpecifics::Marshalling &marshalled) { | |
| llvm::SmallVector<mlir::Type> typeList; | |
| for (auto &typeAndAttr : marshalled) | |
| typeList.emplace_back(std::get<mlir::Type>(typeAndAttr)); | |
| return typeList; | |
| } | |
| /// Taking the address of a function. Modify the signature as needed. | |
| void convertAddrOp(fir::AddrOfOp addrOp) { | |
| rewriter->setInsertionPoint(addrOp); | |
| auto addrTy = mlir::cast<mlir::FunctionType>(addrOp.getType()); | |
| fir::CodeGenSpecifics::Marshalling newInTyAndAttrs; | |
| llvm::SmallVector<mlir::Type> newResTys; | |
| auto loc = addrOp.getLoc(); | |
| for (mlir::Type ty : addrTy.getResults()) { | |
| llvm::TypeSwitch<mlir::Type>(ty) | |
| .Case([&](mlir::ComplexType ty) { | |
| lowerComplexSignatureRes(loc, ty, newResTys, newInTyAndAttrs); | |
| }) | |
| .Case([&](fir::RecordType ty) { | |
| lowerStructSignatureRes(loc, ty, newResTys, newInTyAndAttrs); | |
| }) | |
| .Default([&](mlir::Type ty) { newResTys.push_back(ty); }); | |
| } | |
| llvm::SmallVector<mlir::Type> trailingInTys; | |
| for (mlir::Type ty : addrTy.getInputs()) { | |
| llvm::TypeSwitch<mlir::Type>(ty) | |
| .Case([&](fir::BoxCharType box) { | |
| if (noCharacterConversion) { | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(box)); | |
| } else { | |
| for (auto &tup : specifics->boxcharArgumentType(box.getEleTy())) { | |
| auto attr = std::get<fir::CodeGenSpecifics::Attributes>(tup); | |
| auto argTy = std::get<mlir::Type>(tup); | |
| if (attr.isAppend()) | |
| trailingInTys.push_back(argTy); | |
| else | |
| newInTyAndAttrs.push_back(tup); | |
| } | |
| } | |
| }) | |
| .Case([&](mlir::ComplexType ty) { | |
| lowerComplexSignatureArg(loc, ty, newInTyAndAttrs); | |
| }) | |
| .Case([&](mlir::TupleType tuple) { | |
| if (fir::isCharacterProcedureTuple(tuple)) { | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(tuple.getType(0))); | |
| trailingInTys.push_back(tuple.getType(1)); | |
| } else { | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(ty)); | |
| } | |
| }) | |
| .Case([&](fir::RecordType recTy) { | |
| lowerStructSignatureArg(loc, recTy, newInTyAndAttrs); | |
| }) | |
| .Default([&](mlir::Type ty) { | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(ty)); | |
| }); | |
| } | |
| llvm::SmallVector<mlir::Type> newInTypes = toTypeList(newInTyAndAttrs); | |
| // append trailing input types | |
| newInTypes.insert(newInTypes.end(), trailingInTys.begin(), | |
| trailingInTys.end()); | |
| // replace this op with a new one with the updated signature | |
| auto newTy = rewriter->getFunctionType(newInTypes, newResTys); | |
| auto newOp = fir::AddrOfOp::create(*rewriter, addrOp.getLoc(), newTy, | |
| addrOp.getSymbol()); | |
| replaceOp(addrOp, newOp.getResult()); | |
| } | |
| /// Convert the type signatures on all the functions present in the module. | |
| /// As the type signature is being changed, this must also update the | |
| /// function itself to use any new arguments, etc. | |
| llvm::LogicalResult convertTypes(mlir::ModuleOp mod) { | |
| mlir::MLIRContext *ctx = mod->getContext(); | |
| auto targetCPU = specifics->getTargetCPU(); | |
| mlir::StringAttr targetCPUAttr = | |
| targetCPU.empty() ? nullptr : mlir::StringAttr::get(ctx, targetCPU); | |
| auto tuneCPU = specifics->getTuneCPU(); | |
| mlir::StringAttr tuneCPUAttr = | |
| tuneCPU.empty() ? nullptr : mlir::StringAttr::get(ctx, tuneCPU); | |
| auto targetFeaturesAttr = specifics->getTargetFeatures(); | |
| for (auto fn : mod.getOps<mlir::func::FuncOp>()) { | |
| if (targetCPUAttr) | |
| fn->setAttr("llvm.target_cpu", targetCPUAttr); | |
| if (tuneCPUAttr) | |
| fn->setAttr("llvm.tune_cpu", tuneCPUAttr); | |
| if (targetFeaturesAttr) | |
| fn->setAttr("llvm.target_features", targetFeaturesAttr); | |
| convertSignature<mlir::func::ReturnOp, mlir::func::FuncOp>(fn); | |
| } | |
| for (auto gpuMod : mod.getOps<mlir::gpu::GPUModuleOp>()) { | |
| for (auto fn : gpuMod.getOps<mlir::func::FuncOp>()) | |
| convertSignature<mlir::func::ReturnOp, mlir::func::FuncOp>(fn); | |
| for (auto fn : gpuMod.getOps<mlir::gpu::GPUFuncOp>()) | |
| convertSignature<mlir::gpu::ReturnOp, mlir::gpu::GPUFuncOp>(fn); | |
| } | |
| return mlir::success(); | |
| } | |
| // Returns true if the function should be interoperable with C. | |
| static bool isFuncWithCCallingConvention(mlir::Operation *op) { | |
| auto funcOp = mlir::dyn_cast<mlir::func::FuncOp>(op); | |
| if (!funcOp) | |
| return false; | |
| return op->hasAttrOfType<mlir::UnitAttr>( | |
| fir::FIROpsDialect::getFirRuntimeAttrName()) || | |
| op->hasAttrOfType<mlir::StringAttr>(fir::getSymbolAttrName()); | |
| } | |
| /// If the signature does not need any special target-specific conversions, | |
| /// then it is considered portable for any target, and this function will | |
| /// return `true`. Otherwise, the signature is not portable and `false` is | |
| /// returned. | |
| bool hasPortableSignature(mlir::Type signature, mlir::Operation *op) { | |
| assert(mlir::isa<mlir::FunctionType>(signature)); | |
| auto func = mlir::dyn_cast<mlir::FunctionType>(signature); | |
| bool hasCCallingConv = isFuncWithCCallingConvention(op); | |
| for (auto ty : func.getResults()) | |
| if ((mlir::isa<fir::BoxCharType>(ty) && !noCharacterConversion) || | |
| (fir::isa_complex(ty) && !noComplexConversion) || | |
| (mlir::isa<mlir::IntegerType>(ty) && hasCCallingConv) || | |
| (mlir::isa<fir::RecordType>(ty) && !noStructConversion)) { | |
| LLVM_DEBUG(llvm::dbgs() << "rewrite " << signature << " for target\n"); | |
| return false; | |
| } | |
| for (auto ty : func.getInputs()) | |
| if (((mlir::isa<fir::BoxCharType>(ty) || | |
| fir::isCharacterProcedureTuple(ty)) && | |
| !noCharacterConversion) || | |
| (fir::isa_complex(ty) && !noComplexConversion) || | |
| (mlir::isa<mlir::IntegerType>(ty) && hasCCallingConv) || | |
| (mlir::isa<fir::RecordType>(ty) && !noStructConversion)) { | |
| LLVM_DEBUG(llvm::dbgs() << "rewrite " << signature << " for target\n"); | |
| return false; | |
| } | |
| return true; | |
| } | |
| /// Determine if the signature has host associations. The host association | |
| /// argument may need special target specific rewriting. | |
| template <typename OpTy> | |
| static bool hasHostAssociations(OpTy func) { | |
| std::size_t end = func.getFunctionType().getInputs().size(); | |
| for (std::size_t i = 0; i < end; ++i) | |
| if (func.template getArgAttrOfType<mlir::UnitAttr>( | |
| i, fir::getHostAssocAttrName())) | |
| return true; | |
| return false; | |
| } | |
| /// Rewrite the signatures and body of the `FuncOp`s in the module for | |
| /// the immediately subsequent target code gen. | |
| template <typename ReturnOpTy, typename FuncOpTy> | |
| void convertSignature(FuncOpTy func) { | |
| auto funcTy = mlir::cast<mlir::FunctionType>(func.getFunctionType()); | |
| if (hasPortableSignature(funcTy, func) && !hasHostAssociations(func)) | |
| return; | |
| llvm::SmallVector<mlir::Type> newResTys; | |
| fir::CodeGenSpecifics::Marshalling newInTyAndAttrs; | |
| llvm::SmallVector<std::pair<unsigned, mlir::NamedAttribute>> savedAttrs; | |
| llvm::SmallVector<std::pair<unsigned, mlir::NamedAttribute>> extraAttrs; | |
| llvm::SmallVector<FixupTy> fixups; | |
| llvm::SmallVector<std::pair<unsigned, mlir::NamedAttrList>, 1> resultAttrs; | |
| // Save argument attributes in case there is a shift so we can replace them | |
| // correctly. | |
| for (auto e : llvm::enumerate(funcTy.getInputs())) { | |
| unsigned index = e.index(); | |
| llvm::ArrayRef<mlir::NamedAttribute> attrs = | |
| mlir::function_interface_impl::getArgAttrs(func, index); | |
| for (mlir::NamedAttribute attr : attrs) { | |
| savedAttrs.push_back({index, attr}); | |
| } | |
| } | |
| // Count the number of arguments that have to stay in place at the end of | |
| // the argument list. | |
| unsigned trailingArgs = 0; | |
| if constexpr (std::is_same_v<FuncOpTy, mlir::gpu::GPUFuncOp>) { | |
| trailingArgs = | |
| func.getNumWorkgroupAttributions() + func.getNumPrivateAttributions(); | |
| } | |
| // Convert return value(s) | |
| for (auto ty : funcTy.getResults()) | |
| llvm::TypeSwitch<mlir::Type>(ty) | |
| .Case([&](mlir::ComplexType cmplx) { | |
| if (noComplexConversion) | |
| newResTys.push_back(cmplx); | |
| else | |
| doComplexReturn(func, cmplx, newResTys, newInTyAndAttrs, fixups); | |
| }) | |
| .Case([&](mlir::IntegerType intTy) { | |
| auto m = specifics->integerArgumentType(func.getLoc(), intTy); | |
| assert(m.size() == 1); | |
| auto attr = std::get<fir::CodeGenSpecifics::Attributes>(m[0]); | |
| auto retTy = std::get<mlir::Type>(m[0]); | |
| std::size_t resId = newResTys.size(); | |
| llvm::StringRef extensionAttrName = attr.getIntExtensionAttrName(); | |
| if (!extensionAttrName.empty() && | |
| isFuncWithCCallingConvention(func)) | |
| resultAttrs.emplace_back( | |
| resId, rewriter->getNamedAttr(extensionAttrName, | |
| rewriter->getUnitAttr())); | |
| newResTys.push_back(retTy); | |
| }) | |
| .Case([&](fir::RecordType recTy) { | |
| doStructReturn(func, recTy, newResTys, newInTyAndAttrs, fixups); | |
| }) | |
| .Default([&](mlir::Type ty) { newResTys.push_back(ty); }); | |
| // New index of each original argument, so that saved attributes can follow | |
| // their argument when target expansion shifts arguments to the right. | |
| llvm::SmallVector<unsigned> newArgIndex; | |
| // Convert arguments | |
| llvm::SmallVector<mlir::Type> trailingTys; | |
| for (auto e : llvm::enumerate(funcTy.getInputs())) { | |
| auto ty = e.value(); | |
| unsigned index = e.index(); | |
| newArgIndex.push_back(newInTyAndAttrs.size()); | |
| llvm::TypeSwitch<mlir::Type>(ty) | |
| .Case([&](fir::BoxCharType boxTy) { | |
| if (noCharacterConversion) { | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(boxTy)); | |
| } else { | |
| // Convert a CHARACTER argument type. This can involve separating | |
| // the pointer and the LEN into two arguments and moving the LEN | |
| // argument to the end of the arg list. | |
| for (auto &tup : | |
| specifics->boxcharArgumentType(boxTy.getEleTy())) { | |
| auto attr = std::get<fir::CodeGenSpecifics::Attributes>(tup); | |
| auto argTy = std::get<mlir::Type>(tup); | |
| if (attr.isAppend()) { | |
| trailingTys.push_back(argTy); | |
| } else { | |
| fixups.emplace_back(FixupTy::Codes::Trailing, | |
| newInTyAndAttrs.size(), | |
| trailingTys.size()); | |
| newInTyAndAttrs.push_back(tup); | |
| } | |
| } | |
| } | |
| }) | |
| .Case([&](mlir::ComplexType cmplx) { | |
| doComplexArg(func, cmplx, newInTyAndAttrs, fixups); | |
| }) | |
| .Case([&](mlir::TupleType tuple) { | |
| if (fir::isCharacterProcedureTuple(tuple)) { | |
| fixups.emplace_back(FixupTy::Codes::TrailingCharProc, | |
| newInTyAndAttrs.size(), trailingTys.size()); | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(tuple.getType(0))); | |
| trailingTys.push_back(tuple.getType(1)); | |
| } else { | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(ty)); | |
| } | |
| }) | |
| .Case([&](mlir::IntegerType intTy) { | |
| auto m = specifics->integerArgumentType(func.getLoc(), intTy); | |
| assert(m.size() == 1); | |
| auto attr = std::get<fir::CodeGenSpecifics::Attributes>(m[0]); | |
| auto argNo = newInTyAndAttrs.size(); | |
| llvm::StringRef extensionAttrName = attr.getIntExtensionAttrName(); | |
| if (!extensionAttrName.empty() && | |
| isFuncWithCCallingConvention(func)) | |
| fixups.emplace_back(FixupTy::Codes::ArgumentType, argNo, | |
| [=](FuncOpTy func) { | |
| func.setArgAttr( | |
| argNo, extensionAttrName, | |
| mlir::UnitAttr::get(func.getContext())); | |
| }); | |
| newInTyAndAttrs.push_back(m[0]); | |
| }) | |
| .Case([&](fir::RecordType recTy) { | |
| doStructArg(func, recTy, newInTyAndAttrs, fixups); | |
| }) | |
| .Default([&](mlir::Type ty) { | |
| newInTyAndAttrs.push_back( | |
| fir::CodeGenSpecifics::getTypeAndAttr(ty)); | |
| }); | |
| if (func.template getArgAttrOfType<mlir::UnitAttr>( | |
| index, fir::getHostAssocAttrName())) { | |
| extraAttrs.push_back( | |
| {newInTyAndAttrs.size() - 1, | |
| rewriter->getNamedAttr("llvm.nest", rewriter->getUnitAttr())}); | |
| } | |
| } | |
| // Add the argument at the end if the number of trailing arguments is 0, | |
| // otherwise insert the argument at the appropriate index. | |
| auto addOrInsertArgument = [&](mlir::Type ty, mlir::Location loc) { | |
| unsigned inputIndex = func.front().getArguments().size() - trailingArgs; | |
| auto newArg = trailingArgs == 0 | |
| ? func.front().addArgument(ty, loc) | |
| : func.front().insertArgument(inputIndex, ty, loc); | |
| return newArg; | |
| }; | |
| if (!func.empty()) { | |
| // If the function has a body, then apply the fixups to the arguments and | |
| // return ops as required. These fixups are done in place. | |
| auto loc = func.getLoc(); | |
| const auto fixupSize = fixups.size(); | |
| const auto oldArgTys = func.getFunctionType().getInputs(); | |
| int offset = 0; | |
| for (std::remove_const_t<decltype(fixupSize)> i = 0; i < fixupSize; ++i) { | |
| const auto &fixup = fixups[i]; | |
| mlir::Type fixupType = | |
| fixup.index < newInTyAndAttrs.size() | |
| ? std::get<mlir::Type>(newInTyAndAttrs[fixup.index]) | |
| : mlir::Type{}; | |
| switch (fixup.code) { | |
| case FixupTy::Codes::ArgumentAsLoad: { | |
| // Argument was pass-by-value, but is now pass-by-reference and | |
| // possibly with a different element type. | |
| auto newArg = | |
| func.front().insertArgument(fixup.index, fixupType, loc); | |
| rewriter->setInsertionPointToStart(&func.front()); | |
| auto oldArgTy = | |
| fir::ReferenceType::get(oldArgTys[fixup.index - offset]); | |
| auto cast = fir::ConvertOp::create(*rewriter, loc, oldArgTy, newArg); | |
| auto load = fir::LoadOp::create(*rewriter, loc, cast); | |
| func.getArgument(fixup.index + 1).replaceAllUsesWith(load); | |
| func.front().eraseArgument(fixup.index + 1); | |
| } break; | |
| case FixupTy::Codes::ArgumentType: { | |
| // Argument is pass-by-value, but its type has likely been modified to | |
| // suit the target ABI convention. | |
| auto oldArgTy = oldArgTys[fixup.index - offset]; | |
| // If type did not change, keep the original argument. | |
| if (fixupType == oldArgTy) | |
| break; | |
| auto newArg = | |
| func.front().insertArgument(fixup.index, fixupType, loc); | |
| rewriter->setInsertionPointToStart(&func.front()); | |
| mlir::Value bitcast = convertValueInMemory(loc, newArg, oldArgTy, | |
| /*inputMayBeBigger=*/true); | |
| func.getArgument(fixup.index + 1).replaceAllUsesWith(bitcast); | |
| func.front().eraseArgument(fixup.index + 1); | |
| LLVM_DEBUG(llvm::dbgs() | |
| << "old argument: " << oldArgTy << ", repl: " << bitcast | |
| << ", new argument: " | |
| << func.getArgument(fixup.index).getType() << '\n'); | |
| } break; | |
| case FixupTy::Codes::CharPair: { | |
| // The FIR boxchar argument has been split into a pair of distinct | |
| // arguments that are in juxtaposition to each other. | |
| auto newArg = | |
| func.front().insertArgument(fixup.index, fixupType, loc); | |
| if (fixup.second == 1) { | |
| rewriter->setInsertionPointToStart(&func.front()); | |
| auto boxTy = oldArgTys[fixup.index - offset - fixup.second]; | |
| auto box = fir::EmboxCharOp::create( | |
| *rewriter, loc, boxTy, | |
| func.front().getArgument(fixup.index - 1), newArg); | |
| func.getArgument(fixup.index + 1).replaceAllUsesWith(box); | |
| func.front().eraseArgument(fixup.index + 1); | |
| offset++; | |
| } | |
| } break; | |
| case FixupTy::Codes::ReturnAsStore: { | |
| // The value being returned is now being returned in memory (callee | |
| // stack space) through a hidden reference argument. | |
| auto newArg = | |
| func.front().insertArgument(fixup.index, fixupType, loc); | |
| offset++; | |
| func.walk([&](ReturnOpTy ret) { | |
| rewriter->setInsertionPoint(ret); | |
| auto oldOper = ret.getOperand(0); | |
| auto oldOperTy = fir::ReferenceType::get(oldOper.getType()); | |
| auto cast = | |
| fir::ConvertOp::create(*rewriter, loc, oldOperTy, newArg); | |
| fir::StoreOp::create(*rewriter, loc, oldOper, cast); | |
| ReturnOpTy::create(*rewriter, loc); | |
| ret.erase(); | |
| }); | |
| } break; | |
| case FixupTy::Codes::ReturnType: { | |
| // The function is still returning a value, but its type has likely | |
| // changed to suit the target ABI convention. | |
| func.walk([&](ReturnOpTy ret) { | |
| rewriter->setInsertionPoint(ret); | |
| auto oldOper = ret.getOperand(0); | |
| mlir::Value bitcast = | |
| convertValueInMemory(loc, oldOper, newResTys[fixup.index], | |
| /*inputMayBeBigger=*/false); | |
| ReturnOpTy::create(*rewriter, loc, bitcast); | |
| ret.erase(); | |
| }); | |
| } break; | |
| case FixupTy::Codes::Split: { | |
| // The FIR argument has been split into a pair of distinct arguments | |
| // that are in juxtaposition to each other. (For COMPLEX value or | |
| // derived type passed with VALUE in BIND(C) context). | |
| auto newArg = | |
| func.front().insertArgument(fixup.index, fixupType, loc); | |
| if (fixup.second == 1) { | |
| rewriter->setInsertionPointToStart(&func.front()); | |
| mlir::Value firstArg = func.front().getArgument(fixup.index - 1); | |
| mlir::Type originalTy = | |
| oldArgTys[fixup.index - offset - fixup.second]; | |
| mlir::Type pairTy = originalTy; | |
| if (!fir::isa_complex(originalTy)) { | |
| pairTy = mlir::TupleType::get( | |
| originalTy.getContext(), | |
| mlir::TypeRange{firstArg.getType(), newArg.getType()}); | |
| } | |
| auto undef = fir::UndefOp::create(*rewriter, loc, pairTy); | |
| auto iTy = rewriter->getIntegerType(32); | |
| auto zero = rewriter->getIntegerAttr(iTy, 0); | |
| auto one = rewriter->getIntegerAttr(iTy, 1); | |
| mlir::Value pair1 = fir::InsertValueOp::create( | |
| *rewriter, loc, pairTy, undef, firstArg, | |
| rewriter->getArrayAttr(zero)); | |
| mlir::Value pair = | |
| fir::InsertValueOp::create(*rewriter, loc, pairTy, pair1, | |
| newArg, rewriter->getArrayAttr(one)); | |
| // Cast local argument tuple to original type via memory if needed. | |
| if (pairTy != originalTy) | |
| pair = convertValueInMemory(loc, pair, originalTy, | |
| /*inputMayBeBigger=*/true); | |
| func.getArgument(fixup.index + 1).replaceAllUsesWith(pair); | |
| func.front().eraseArgument(fixup.index + 1); | |
| offset++; | |
| } | |
| } break; | |
| case FixupTy::Codes::Trailing: { | |
| // The FIR argument has been split into a pair of distinct arguments. | |
| // The first part of the pair appears in the original argument | |
| // position. The second part of the pair is appended after all the | |
| // original arguments. (Boxchar arguments.) | |
| auto newBufArg = | |
| func.front().insertArgument(fixup.index, fixupType, loc); | |
| auto newLenArg = addOrInsertArgument(trailingTys[fixup.second], loc); | |
| auto boxTy = oldArgTys[fixup.index - offset]; | |
| rewriter->setInsertionPointToStart(&func.front()); | |
| auto box = fir::EmboxCharOp::create(*rewriter, loc, boxTy, newBufArg, | |
| newLenArg); | |
| func.getArgument(fixup.index + 1).replaceAllUsesWith(box); | |
| func.front().eraseArgument(fixup.index + 1); | |
| } break; | |
| case FixupTy::Codes::TrailingCharProc: { | |
| // The FIR character procedure argument tuple must be split into a | |
| // pair of distinct arguments. The first part of the pair appears in | |
| // the original argument position. The second part of the pair is | |
| // appended after all the original arguments. | |
| auto newProcPointerArg = | |
| func.front().insertArgument(fixup.index, fixupType, loc); | |
| auto newLenArg = addOrInsertArgument(trailingTys[fixup.second], loc); | |
| auto tupleType = oldArgTys[fixup.index - offset]; | |
| rewriter->setInsertionPointToStart(&func.front()); | |
| fir::FirOpBuilder builder(*rewriter, getModule()); | |
| auto tuple = fir::factory::createCharacterProcedureTuple( | |
| builder, loc, tupleType, newProcPointerArg, newLenArg); | |
| func.getArgument(fixup.index + 1).replaceAllUsesWith(tuple); | |
| func.front().eraseArgument(fixup.index + 1); | |
| } break; | |
| } | |
| } | |
| } | |
| llvm::SmallVector<mlir::Type> newInTypes = toTypeList(newInTyAndAttrs); | |
| // Set the new type and finalize the arguments, etc. | |
| newInTypes.insert(newInTypes.end(), trailingTys.begin(), trailingTys.end()); | |
| auto newFuncTy = | |
| mlir::FunctionType::get(func.getContext(), newInTypes, newResTys); | |
| LLVM_DEBUG(llvm::dbgs() << "new func: " << newFuncTy << '\n'); | |
| func.setType(newFuncTy); | |
| for (std::pair<unsigned, mlir::NamedAttribute> extraAttr : extraAttrs) | |
| func.setArgAttr(extraAttr.first, extraAttr.second.getName(), | |
| extraAttr.second.getValue()); | |
| for (auto [resId, resAttrList] : resultAttrs) | |
| for (mlir::NamedAttribute resAttr : resAttrList) | |
| func.setResultAttr(resId, resAttr.getName(), resAttr.getValue()); | |
| // Move each saved argument attribute to the index its argument now | |
| // occupies. Target expansion (e.g. splitting a complex value into two | |
| // scalars) can insert arguments mid-list and shift the following ones; an | |
| // attribute such as fir.host_assoc must follow its argument rather than | |
| // stay on the old index and land on an unrelated one. This keeps the | |
| // attributes consistent (it may not cause a miscompile today). Remove all | |
| // before setting so a moved attribute does not clobber another argument's | |
| // attribute. | |
| for (std::pair<unsigned, mlir::NamedAttribute> savedAttr : savedAttrs) | |
| func.removeArgAttr(savedAttr.first, savedAttr.second.getName()); | |
| for (std::pair<unsigned, mlir::NamedAttribute> savedAttr : savedAttrs) | |
| func.setArgAttr(newArgIndex[savedAttr.first], savedAttr.second.getName(), | |
| savedAttr.second.getValue()); | |
| for (auto &fixup : fixups) { | |
| if constexpr (std::is_same_v<FuncOpTy, mlir::func::FuncOp>) | |
| if (fixup.finalizer) | |
| (*fixup.finalizer)(func); | |
| if constexpr (std::is_same_v<FuncOpTy, mlir::gpu::GPUFuncOp>) | |
| if (fixup.gpuFinalizer) | |
| (*fixup.gpuFinalizer)(func); | |
| } | |
| } | |
| template <typename OpTy, typename Ty, typename FIXUPS> | |
| void doReturn(OpTy func, Ty &newResTys, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, | |
| FIXUPS &fixups, fir::CodeGenSpecifics::Marshalling &m) { | |
| assert(m.size() == 1 && | |
| "expect result to be turned into single argument or result so far"); | |
| auto &tup = m[0]; | |
| auto attr = std::get<fir::CodeGenSpecifics::Attributes>(tup); | |
| auto argTy = std::get<mlir::Type>(tup); | |
| if (attr.isSRet()) { | |
| unsigned argNo = newInTyAndAttrs.size(); | |
| if (auto align = attr.getAlignment()) | |
| fixups.emplace_back( | |
| FixupTy::Codes::ReturnAsStore, argNo, [=](OpTy func) { | |
| auto elemType = fir::dyn_cast_ptrOrBoxEleTy( | |
| func.getFunctionType().getInput(argNo)); | |
| func.setArgAttr(argNo, "llvm.sret", | |
| mlir::TypeAttr::get(elemType)); | |
| func.setArgAttr(argNo, "llvm.align", | |
| rewriter->getIntegerAttr( | |
| rewriter->getIntegerType(32), align)); | |
| }); | |
| else | |
| fixups.emplace_back(FixupTy::Codes::ReturnAsStore, argNo, | |
| [=](OpTy func) { | |
| auto elemType = fir::dyn_cast_ptrOrBoxEleTy( | |
| func.getFunctionType().getInput(argNo)); | |
| func.setArgAttr(argNo, "llvm.sret", | |
| mlir::TypeAttr::get(elemType)); | |
| }); | |
| newInTyAndAttrs.push_back(tup); | |
| return; | |
| } | |
| if (auto align = attr.getAlignment()) | |
| fixups.emplace_back( | |
| FixupTy::Codes::ReturnType, newResTys.size(), [=](OpTy func) { | |
| func.setArgAttr( | |
| newResTys.size(), "llvm.align", | |
| rewriter->getIntegerAttr(rewriter->getIntegerType(32), align)); | |
| }); | |
| else | |
| fixups.emplace_back(FixupTy::Codes::ReturnType, newResTys.size()); | |
| newResTys.push_back(argTy); | |
| } | |
| /// Convert a complex return value. This can involve converting the return | |
| /// value to a "hidden" first argument or packing the complex into a wide | |
| /// GPR. | |
| template <typename OpTy, typename Ty, typename FIXUPS> | |
| void doComplexReturn(OpTy func, mlir::ComplexType cmplx, Ty &newResTys, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, | |
| FIXUPS &fixups) { | |
| if (noComplexConversion) { | |
| newResTys.push_back(cmplx); | |
| return; | |
| } | |
| auto m = | |
| specifics->complexReturnType(func.getLoc(), cmplx.getElementType()); | |
| doReturn(func, newResTys, newInTyAndAttrs, fixups, m); | |
| } | |
| template <typename OpTy, typename Ty, typename FIXUPS> | |
| void doStructReturn(OpTy func, fir::RecordType recTy, Ty &newResTys, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, | |
| FIXUPS &fixups) { | |
| if (noStructConversion) { | |
| newResTys.push_back(recTy); | |
| return; | |
| } | |
| auto m = specifics->structReturnType(func.getLoc(), recTy); | |
| doReturn(func, newResTys, newInTyAndAttrs, fixups, m); | |
| } | |
| template <typename OpTy, typename FIXUPS> | |
| void createFuncOpArgFixups( | |
| OpTy func, fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, | |
| fir::CodeGenSpecifics::Marshalling &argsInTys, FIXUPS &fixups) { | |
| const auto fixupCode = argsInTys.size() > 1 ? FixupTy::Codes::Split | |
| : FixupTy::Codes::ArgumentType; | |
| for (auto e : llvm::enumerate(argsInTys)) { | |
| auto &tup = e.value(); | |
| auto index = e.index(); | |
| auto attr = std::get<fir::CodeGenSpecifics::Attributes>(tup); | |
| auto argNo = newInTyAndAttrs.size(); | |
| if (attr.isByVal() || attr.isIndirect()) { | |
| // In both cases the callee receives a pointer to a copy made by the | |
| // caller and loads the value from it. They differ only in the attribute | |
| // attached to the argument: `byval` marks it as a by-value aggregate | |
| // for the target to lower, whereas an indirect argument is passed as a | |
| // plain pointer. | |
| const bool setByVal = attr.isByVal(); | |
| const unsigned align = attr.getAlignment(); | |
| fixups.emplace_back(FixupTy::Codes::ArgumentAsLoad, argNo, | |
| [=](OpTy func) { | |
| if (setByVal) { | |
| auto elemType = fir::dyn_cast_ptrOrBoxEleTy( | |
| func.getFunctionType().getInput(argNo)); | |
| func.setArgAttr(argNo, "llvm.byval", | |
| mlir::TypeAttr::get(elemType)); | |
| } | |
| if (align) | |
| func.setArgAttr( | |
| argNo, "llvm.align", | |
| rewriter->getIntegerAttr( | |
| rewriter->getIntegerType(32), align)); | |
| }); | |
| } else { | |
| if (auto align = attr.getAlignment()) | |
| fixups.emplace_back( | |
| fixupCode, argNo, index, [=](OpTy func) { | |
| func.setArgAttr(argNo, "llvm.align", | |
| rewriter->getIntegerAttr( | |
| rewriter->getIntegerType(32), align)); | |
| }); | |
| else | |
| fixups.emplace_back(fixupCode, argNo, index); | |
| } | |
| newInTyAndAttrs.push_back(tup); | |
| } | |
| } | |
| /// Convert a complex argument value. This can involve storing the value to | |
| /// a temporary memory location or factoring the value into two distinct | |
| /// arguments. | |
| template <typename OpTy, typename FIXUPS> | |
| void doComplexArg(OpTy func, mlir::ComplexType cmplx, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, | |
| FIXUPS &fixups) { | |
| if (noComplexConversion) { | |
| newInTyAndAttrs.push_back(fir::CodeGenSpecifics::getTypeAndAttr(cmplx)); | |
| return; | |
| } | |
| auto cplxArgs = | |
| specifics->complexArgumentType(func.getLoc(), cmplx.getElementType()); | |
| createFuncOpArgFixups(func, newInTyAndAttrs, cplxArgs, fixups); | |
| } | |
| template <typename OpTy, typename FIXUPS> | |
| void doStructArg(OpTy func, fir::RecordType recTy, | |
| fir::CodeGenSpecifics::Marshalling &newInTyAndAttrs, | |
| FIXUPS &fixups) { | |
| if (noStructConversion) { | |
| newInTyAndAttrs.push_back(fir::CodeGenSpecifics::getTypeAndAttr(recTy)); | |
| return; | |
| } | |
| auto structArgs = | |
| specifics->structArgumentType(func.getLoc(), recTy, newInTyAndAttrs); | |
| createFuncOpArgFixups(func, newInTyAndAttrs, structArgs, fixups); | |
| } | |
| private: | |
| // Replace `op` and remove it. | |
| void replaceOp(mlir::Operation *op, mlir::ValueRange newValues) { | |
| llvm::SmallVector<mlir::Value> casts; | |
| for (auto [oldValue, newValue] : llvm::zip(op->getResults(), newValues)) { | |
| if (oldValue.getType() == newValue.getType()) | |
| casts.push_back(newValue); | |
| else | |
| casts.push_back(fir::ConvertOp::create(*rewriter, op->getLoc(), | |
| oldValue.getType(), newValue)); | |
| } | |
| op->replaceAllUsesWith(casts); | |
| op->dropAllReferences(); | |
| op->erase(); | |
| } | |
| inline void setMembers(fir::CodeGenSpecifics *s, mlir::OpBuilder *r, | |
| mlir::DataLayout *dl) { | |
| specifics = s; | |
| rewriter = r; | |
| dataLayout = dl; | |
| } | |
| inline void clearMembers() { setMembers(nullptr, nullptr, nullptr); } | |
| // Inserts a call to llvm.stacksave at the current insertion | |
| // point and the given location. Returns the call's result Value. | |
| inline mlir::Value genStackSave(mlir::Location loc) { | |
| fir::FirOpBuilder builder(*rewriter, getModule()); | |
| return builder.genStackSave(loc); | |
| } | |
| // Inserts a call to llvm.stackrestore at the current insertion | |
| // point and the given location and argument. | |
| inline void genStackRestore(mlir::Location loc, mlir::Value sp) { | |
| fir::FirOpBuilder builder(*rewriter, getModule()); | |
| return builder.genStackRestore(loc, sp); | |
| } | |
| fir::CodeGenSpecifics *specifics = nullptr; | |
| mlir::OpBuilder *rewriter = nullptr; | |
| mlir::DataLayout *dataLayout = nullptr; | |
| }; | |
| } // namespace |