4451 lines · cpp
1//===-- CodeGen.cpp -- bridge to lower to LLVM ----------------------------===//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// Coding style: https://mlir.llvm.org/getting_started/DeveloperGuide/10//11//===----------------------------------------------------------------------===//12 13#include "flang/Optimizer/CodeGen/CodeGen.h"14 15#include "flang/Optimizer/Builder/CUFCommon.h"16#include "flang/Optimizer/CodeGen/CodeGenOpenMP.h"17#include "flang/Optimizer/CodeGen/FIROpPatterns.h"18#include "flang/Optimizer/CodeGen/LLVMInsertChainFolder.h"19#include "flang/Optimizer/CodeGen/TypeConverter.h"20#include "flang/Optimizer/Dialect/FIRAttr.h"21#include "flang/Optimizer/Dialect/FIRCG/CGOps.h"22#include "flang/Optimizer/Dialect/FIRDialect.h"23#include "flang/Optimizer/Dialect/FIROps.h"24#include "flang/Optimizer/Dialect/FIRType.h"25#include "flang/Optimizer/Support/DataLayout.h"26#include "flang/Optimizer/Support/InternalNames.h"27#include "flang/Optimizer/Support/TypeCode.h"28#include "flang/Optimizer/Support/Utils.h"29#include "flang/Runtime/CUDA/descriptor.h"30#include "flang/Runtime/CUDA/memory.h"31#include "flang/Runtime/allocator-registry-consts.h"32#include "flang/Runtime/descriptor-consts.h"33#include "flang/Semantics/runtime-type-info.h"34#include "mlir/Conversion/ArithCommon/AttrToLLVMConverter.h"35#include "mlir/Conversion/ArithToLLVM/ArithToLLVM.h"36#include "mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h"37#include "mlir/Conversion/ComplexToROCDLLibraryCalls/ComplexToROCDLLibraryCalls.h"38#include "mlir/Conversion/ComplexToStandard/ComplexToStandard.h"39#include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h"40#include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h"41#include "mlir/Conversion/IndexToLLVM/IndexToLLVM.h"42#include "mlir/Conversion/LLVMCommon/Pattern.h"43#include "mlir/Conversion/MathToFuncs/MathToFuncs.h"44#include "mlir/Conversion/MathToLLVM/MathToLLVM.h"45#include "mlir/Conversion/MathToLibm/MathToLibm.h"46#include "mlir/Conversion/MathToROCDL/MathToROCDL.h"47#include "mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h"48#include "mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h"49#include "mlir/Dialect/Arith/IR/Arith.h"50#include "mlir/Dialect/DLTI/DLTI.h"51#include "mlir/Dialect/GPU/IR/GPUDialect.h"52#include "mlir/Dialect/LLVMIR/LLVMAttrs.h"53#include "mlir/Dialect/LLVMIR/LLVMDialect.h"54#include "mlir/Dialect/LLVMIR/NVVMDialect.h"55#include "mlir/Dialect/LLVMIR/Transforms/AddComdats.h"56#include "mlir/Dialect/OpenACC/OpenACC.h"57#include "mlir/Dialect/OpenMP/OpenMPDialect.h"58#include "mlir/IR/BuiltinTypes.h"59#include "mlir/IR/Matchers.h"60#include "mlir/Pass/Pass.h"61#include "mlir/Pass/PassManager.h"62#include "mlir/Target/LLVMIR/Import.h"63#include "mlir/Target/LLVMIR/ModuleTranslation.h"64#include "llvm/ADT/ArrayRef.h"65#include "llvm/ADT/TypeSwitch.h"66 67namespace fir {68#define GEN_PASS_DEF_FIRTOLLVMLOWERING69#include "flang/Optimizer/CodeGen/CGPasses.h.inc"70} // namespace fir71 72#define DEBUG_TYPE "flang-codegen"73 74// TODO: This should really be recovered from the specified target.75static constexpr unsigned defaultAlign = 8;76 77/// `fir.box` attribute values as defined for CFI_attribute_t in78/// flang/ISO_Fortran_binding.h.79static constexpr unsigned kAttrPointer = CFI_attribute_pointer;80static constexpr unsigned kAttrAllocatable = CFI_attribute_allocatable;81 82static inline mlir::Type getLlvmPtrType(mlir::MLIRContext *context,83 unsigned addressSpace = 0) {84 return mlir::LLVM::LLVMPointerType::get(context, addressSpace);85}86 87static inline mlir::Type getI8Type(mlir::MLIRContext *context) {88 return mlir::IntegerType::get(context, 8);89}90 91static mlir::Block *createBlock(mlir::ConversionPatternRewriter &rewriter,92 mlir::Block *insertBefore) {93 assert(insertBefore && "expected valid insertion block");94 return rewriter.createBlock(insertBefore->getParent(),95 mlir::Region::iterator(insertBefore));96}97 98/// Extract constant from a value that must be the result of one of the99/// ConstantOp operations.100static int64_t getConstantIntValue(mlir::Value val) {101 if (auto constVal = fir::getIntIfConstant(val))102 return *constVal;103 fir::emitFatalError(val.getLoc(), "must be a constant");104}105 106static unsigned getTypeDescFieldId(mlir::Type ty) {107 auto isArray = mlir::isa<fir::SequenceType>(fir::dyn_cast_ptrOrBoxEleTy(ty));108 return isArray ? kOptTypePtrPosInBox : kDimsPosInBox;109}110static unsigned getLenParamFieldId(mlir::Type ty) {111 return getTypeDescFieldId(ty) + 1;112}113 114static llvm::SmallVector<mlir::NamedAttribute>115addLLVMOpBundleAttrs(mlir::ConversionPatternRewriter &rewriter,116 llvm::ArrayRef<mlir::NamedAttribute> attrs,117 int32_t numCallOperands) {118 llvm::SmallVector<mlir::NamedAttribute> newAttrs;119 newAttrs.reserve(attrs.size() + 2);120 121 for (mlir::NamedAttribute attr : attrs) {122 if (attr.getName() != "operandSegmentSizes")123 newAttrs.push_back(attr);124 }125 126 newAttrs.push_back(rewriter.getNamedAttr(127 "operandSegmentSizes",128 rewriter.getDenseI32ArrayAttr({numCallOperands, 0})));129 newAttrs.push_back(rewriter.getNamedAttr("op_bundle_sizes",130 rewriter.getDenseI32ArrayAttr({})));131 return newAttrs;132}133 134namespace {135 136// Replaces an existing operation with an AddressOfOp or an AddrSpaceCastOp137// depending on the existing address spaces of the type.138mlir::Value replaceWithAddrOfOrASCast(mlir::ConversionPatternRewriter &rewriter,139 mlir::Location loc,140 std::uint64_t globalAS,141 std::uint64_t programAS,142 llvm::StringRef symName, mlir::Type type,143 mlir::Operation *replaceOp = nullptr) {144 if (mlir::isa<mlir::LLVM::LLVMPointerType>(type)) {145 if (globalAS != programAS) {146 auto llvmAddrOp = mlir::LLVM::AddressOfOp::create(147 rewriter, loc, getLlvmPtrType(rewriter.getContext(), globalAS),148 symName);149 if (replaceOp)150 return rewriter.replaceOpWithNewOp<mlir::LLVM::AddrSpaceCastOp>(151 replaceOp, ::getLlvmPtrType(rewriter.getContext(), programAS),152 llvmAddrOp);153 return mlir::LLVM::AddrSpaceCastOp::create(154 rewriter, loc, getLlvmPtrType(rewriter.getContext(), programAS),155 llvmAddrOp);156 }157 158 if (replaceOp)159 return rewriter.replaceOpWithNewOp<mlir::LLVM::AddressOfOp>(160 replaceOp, getLlvmPtrType(rewriter.getContext(), globalAS), symName);161 return mlir::LLVM::AddressOfOp::create(162 rewriter, loc, getLlvmPtrType(rewriter.getContext(), globalAS),163 symName);164 }165 166 if (replaceOp)167 return rewriter.replaceOpWithNewOp<mlir::LLVM::AddressOfOp>(replaceOp, type,168 symName);169 return mlir::LLVM::AddressOfOp::create(rewriter, loc, type, symName);170}171 172/// Lower `fir.address_of` operation to `llvm.address_of` operation.173struct AddrOfOpConversion : public fir::FIROpConversion<fir::AddrOfOp> {174 using FIROpConversion::FIROpConversion;175 176 llvm::LogicalResult177 matchAndRewrite(fir::AddrOfOp addr, OpAdaptor adaptor,178 mlir::ConversionPatternRewriter &rewriter) const override {179 180 if (auto gpuMod = addr->getParentOfType<mlir::gpu::GPUModuleOp>()) {181 auto global = gpuMod.lookupSymbol<mlir::LLVM::GlobalOp>(addr.getSymbol());182 replaceWithAddrOfOrASCast(183 rewriter, addr->getLoc(),184 global ? global.getAddrSpace() : getGlobalAddressSpace(rewriter),185 getProgramAddressSpace(rewriter),186 global ? global.getSymName()187 : addr.getSymbol().getRootReference().getValue(),188 convertType(addr.getType()), addr);189 return mlir::success();190 }191 192 auto global = addr->getParentOfType<mlir::ModuleOp>()193 .lookupSymbol<mlir::LLVM::GlobalOp>(addr.getSymbol());194 replaceWithAddrOfOrASCast(195 rewriter, addr->getLoc(),196 global ? global.getAddrSpace() : getGlobalAddressSpace(rewriter),197 getProgramAddressSpace(rewriter),198 global ? global.getSymName()199 : addr.getSymbol().getRootReference().getValue(),200 convertType(addr.getType()), addr);201 return mlir::success();202 }203};204} // namespace205 206/// Lookup the function to compute the memory size of this parametric derived207/// type. The size of the object may depend on the LEN type parameters of the208/// derived type.209static mlir::LLVM::LLVMFuncOp210getDependentTypeMemSizeFn(fir::RecordType recTy, fir::AllocaOp op,211 mlir::ConversionPatternRewriter &rewriter) {212 auto module = op->getParentOfType<mlir::ModuleOp>();213 std::string name = recTy.getName().str() + "P.mem.size";214 if (auto memSizeFunc = module.lookupSymbol<mlir::LLVM::LLVMFuncOp>(name))215 return memSizeFunc;216 TODO(op.getLoc(), "did not find allocation function");217}218 219namespace {220struct DeclareOpConversion : public fir::FIROpConversion<fir::cg::XDeclareOp> {221public:222 using FIROpConversion::FIROpConversion;223 llvm::LogicalResult224 matchAndRewrite(fir::cg::XDeclareOp declareOp, OpAdaptor adaptor,225 mlir::ConversionPatternRewriter &rewriter) const override {226 auto memRef = adaptor.getOperands()[0];227 if (auto fusedLoc = mlir::dyn_cast<mlir::FusedLoc>(declareOp.getLoc())) {228 if (auto varAttr =229 mlir::dyn_cast_or_null<mlir::LLVM::DILocalVariableAttr>(230 fusedLoc.getMetadata())) {231 mlir::LLVM::DbgDeclareOp::create(rewriter, memRef.getLoc(), memRef,232 varAttr, nullptr);233 }234 }235 rewriter.replaceOp(declareOp, memRef);236 return mlir::success();237 }238};239} // namespace240 241namespace {242/// convert to LLVM IR dialect `alloca`243struct AllocaOpConversion : public fir::FIROpConversion<fir::AllocaOp> {244 using FIROpConversion::FIROpConversion;245 246 llvm::LogicalResult247 matchAndRewrite(fir::AllocaOp alloc, OpAdaptor adaptor,248 mlir::ConversionPatternRewriter &rewriter) const override {249 mlir::ValueRange operands = adaptor.getOperands();250 auto loc = alloc.getLoc();251 mlir::Type ity = lowerTy().indexType();252 unsigned i = 0;253 mlir::Value size = fir::genConstantIndex(loc, ity, rewriter, 1).getResult();254 mlir::Type firObjType = fir::unwrapRefType(alloc.getType());255 mlir::Type llvmObjectType = convertObjectType(firObjType);256 if (alloc.hasLenParams()) {257 unsigned end = alloc.numLenParams();258 llvm::SmallVector<mlir::Value> lenParams;259 for (; i < end; ++i)260 lenParams.push_back(operands[i]);261 mlir::Type scalarType = fir::unwrapSequenceType(alloc.getInType());262 if (auto chrTy = mlir::dyn_cast<fir::CharacterType>(scalarType)) {263 fir::CharacterType rawCharTy = fir::CharacterType::getUnknownLen(264 chrTy.getContext(), chrTy.getFKind());265 llvmObjectType = convertType(rawCharTy);266 assert(end == 1);267 size = integerCast(loc, rewriter, ity, lenParams[0], /*fold=*/true);268 } else if (auto recTy = mlir::dyn_cast<fir::RecordType>(scalarType)) {269 mlir::LLVM::LLVMFuncOp memSizeFn =270 getDependentTypeMemSizeFn(recTy, alloc, rewriter);271 if (!memSizeFn)272 emitError(loc, "did not find allocation function");273 mlir::NamedAttribute attr = rewriter.getNamedAttr(274 "callee", mlir::SymbolRefAttr::get(memSizeFn));275 auto call = mlir::LLVM::CallOp::create(276 rewriter, loc, ity, lenParams,277 addLLVMOpBundleAttrs(rewriter, {attr}, lenParams.size()));278 size = call.getResult();279 llvmObjectType = ::getI8Type(alloc.getContext());280 } else {281 return emitError(loc, "unexpected type ")282 << scalarType << " with type parameters";283 }284 }285 if (auto scaleSize = fir::genAllocationScaleSize(286 alloc.getLoc(), alloc.getInType(), ity, rewriter))287 size =288 rewriter.createOrFold<mlir::LLVM::MulOp>(loc, ity, size, scaleSize);289 if (alloc.hasShapeOperands()) {290 unsigned end = operands.size();291 for (; i < end; ++i)292 size = rewriter.createOrFold<mlir::LLVM::MulOp>(293 loc, ity, size,294 integerCast(loc, rewriter, ity, operands[i], /*fold=*/true));295 }296 297 unsigned allocaAs = getAllocaAddressSpace(rewriter);298 unsigned programAs = getProgramAddressSpace(rewriter);299 300 if (mlir::isa<mlir::LLVM::ConstantOp>(size.getDefiningOp())) {301 // Set the Block in which the llvm alloca should be inserted.302 mlir::Operation *parentOp = rewriter.getInsertionBlock()->getParentOp();303 mlir::Region *parentRegion = rewriter.getInsertionBlock()->getParent();304 mlir::Block *insertBlock =305 getBlockForAllocaInsert(parentOp, parentRegion);306 307 // The old size might have had multiple users, some at a broader scope308 // than we can safely outline the alloca to. As it is only an309 // llvm.constant operation, it is faster to clone it than to calculate the310 // dominance to see if it really should be moved.311 mlir::Operation *clonedSize = rewriter.clone(*size.getDefiningOp());312 size = clonedSize->getResult(0);313 clonedSize->moveBefore(&insertBlock->front());314 rewriter.setInsertionPointAfter(size.getDefiningOp());315 }316 317 // NOTE: we used to pass alloc->getAttrs() in the builder for non opaque318 // pointers! Only propagate pinned and bindc_name to help debugging, but319 // this should have no functional purpose (and passing the operand segment320 // attribute like before is certainly bad).321 auto llvmAlloc = mlir::LLVM::AllocaOp::create(322 rewriter, loc, ::getLlvmPtrType(alloc.getContext(), allocaAs),323 llvmObjectType, size);324 if (alloc.getPinned())325 llvmAlloc->setDiscardableAttr(alloc.getPinnedAttrName(),326 alloc.getPinnedAttr());327 if (alloc.getBindcName())328 llvmAlloc->setDiscardableAttr(alloc.getBindcNameAttrName(),329 alloc.getBindcNameAttr());330 if (allocaAs == programAs) {331 rewriter.replaceOp(alloc, llvmAlloc);332 } else {333 // if our allocation address space, is not the same as the program address334 // space, then we must emit a cast to the program address space before335 // use. An example case would be on AMDGPU, where the allocation address336 // space is the numeric value 5 (private), and the program address space337 // is 0 (generic).338 rewriter.replaceOpWithNewOp<mlir::LLVM::AddrSpaceCastOp>(339 alloc, ::getLlvmPtrType(alloc.getContext(), programAs), llvmAlloc);340 }341 342 return mlir::success();343 }344};345} // namespace346 347namespace {348 349static bool isInGlobalOp(mlir::ConversionPatternRewriter &rewriter) {350 auto *thisBlock = rewriter.getInsertionBlock();351 return thisBlock && mlir::isa<mlir::LLVM::GlobalOp>(thisBlock->getParentOp());352}353 354// Inside a fir.global, the input box was produced as an llvm.struct<>355// because objects cannot be handled in memory inside a fir.global body that356// must be constant foldable. However, the type translation are not357// contextual, so the fir.box<T> type of the operation that produced the358// fir.box was translated to an llvm.ptr<llvm.struct<>> and the MLIR pass359// manager inserted a builtin.unrealized_conversion_cast that was inserted360// and needs to be removed here.361// This should be called by any pattern operating on operations that are362// accepting fir.box inputs and are used in fir.global.363static mlir::Value364fixBoxInputInsideGlobalOp(mlir::ConversionPatternRewriter &rewriter,365 mlir::Value box) {366 if (isInGlobalOp(rewriter))367 if (auto unrealizedCast =368 box.getDefiningOp<mlir::UnrealizedConversionCastOp>())369 return unrealizedCast.getInputs()[0];370 return box;371}372 373/// Lower `fir.box_addr` to the sequence of operations to extract the first374/// element of the box.375struct BoxAddrOpConversion : public fir::FIROpConversion<fir::BoxAddrOp> {376 using FIROpConversion::FIROpConversion;377 378 llvm::LogicalResult379 matchAndRewrite(fir::BoxAddrOp boxaddr, OpAdaptor adaptor,380 mlir::ConversionPatternRewriter &rewriter) const override {381 mlir::Value a = adaptor.getOperands()[0];382 auto loc = boxaddr.getLoc();383 if (auto argty =384 mlir::dyn_cast<fir::BaseBoxType>(boxaddr.getVal().getType())) {385 a = fixBoxInputInsideGlobalOp(rewriter, a);386 TypePair boxTyPair = getBoxTypePair(argty);387 rewriter.replaceOp(boxaddr,388 getBaseAddrFromBox(loc, boxTyPair, a, rewriter));389 } else {390 rewriter.replaceOpWithNewOp<mlir::LLVM::ExtractValueOp>(boxaddr, a, 0);391 }392 return mlir::success();393 }394};395 396/// Convert `!fir.boxchar_len` to `!llvm.extractvalue` for the 2nd part of the397/// boxchar.398struct BoxCharLenOpConversion : public fir::FIROpConversion<fir::BoxCharLenOp> {399 using FIROpConversion::FIROpConversion;400 401 llvm::LogicalResult402 matchAndRewrite(fir::BoxCharLenOp boxCharLen, OpAdaptor adaptor,403 mlir::ConversionPatternRewriter &rewriter) const override {404 mlir::Value boxChar = adaptor.getOperands()[0];405 mlir::Location loc = boxChar.getLoc();406 mlir::Type returnValTy = boxCharLen.getResult().getType();407 408 constexpr int boxcharLenIdx = 1;409 auto len = mlir::LLVM::ExtractValueOp::create(rewriter, loc, boxChar,410 boxcharLenIdx);411 mlir::Value lenAfterCast = integerCast(loc, rewriter, returnValTy, len);412 rewriter.replaceOp(boxCharLen, lenAfterCast);413 414 return mlir::success();415 }416};417 418/// Lower `fir.box_dims` to a sequence of operations to extract the requested419/// dimension information from the boxed value.420/// Result in a triple set of GEPs and loads.421struct BoxDimsOpConversion : public fir::FIROpConversion<fir::BoxDimsOp> {422 using FIROpConversion::FIROpConversion;423 424 llvm::LogicalResult425 matchAndRewrite(fir::BoxDimsOp boxdims, OpAdaptor adaptor,426 mlir::ConversionPatternRewriter &rewriter) const override {427 llvm::SmallVector<mlir::Type, 3> resultTypes = {428 convertType(boxdims.getResult(0).getType()),429 convertType(boxdims.getResult(1).getType()),430 convertType(boxdims.getResult(2).getType()),431 };432 TypePair boxTyPair = getBoxTypePair(boxdims.getVal().getType());433 auto results = getDimsFromBox(boxdims.getLoc(), resultTypes, boxTyPair,434 adaptor.getOperands()[0],435 adaptor.getOperands()[1], rewriter);436 rewriter.replaceOp(boxdims, results);437 return mlir::success();438 }439};440 441/// Lower `fir.box_elesize` to a sequence of operations ro extract the size of442/// an element in the boxed value.443struct BoxEleSizeOpConversion : public fir::FIROpConversion<fir::BoxEleSizeOp> {444 using FIROpConversion::FIROpConversion;445 446 llvm::LogicalResult447 matchAndRewrite(fir::BoxEleSizeOp boxelesz, OpAdaptor adaptor,448 mlir::ConversionPatternRewriter &rewriter) const override {449 mlir::Value box = adaptor.getOperands()[0];450 auto loc = boxelesz.getLoc();451 auto ty = convertType(boxelesz.getType());452 TypePair boxTyPair = getBoxTypePair(boxelesz.getVal().getType());453 auto elemSize = getElementSizeFromBox(loc, ty, boxTyPair, box, rewriter);454 rewriter.replaceOp(boxelesz, elemSize);455 return mlir::success();456 }457};458 459/// Lower `fir.box_isalloc` to a sequence of operations to determine if the460/// boxed value was from an ALLOCATABLE entity.461struct BoxIsAllocOpConversion : public fir::FIROpConversion<fir::BoxIsAllocOp> {462 using FIROpConversion::FIROpConversion;463 464 llvm::LogicalResult465 matchAndRewrite(fir::BoxIsAllocOp boxisalloc, OpAdaptor adaptor,466 mlir::ConversionPatternRewriter &rewriter) const override {467 mlir::Value box = adaptor.getOperands()[0];468 auto loc = boxisalloc.getLoc();469 TypePair boxTyPair = getBoxTypePair(boxisalloc.getVal().getType());470 mlir::Value check =471 genBoxAttributeCheck(loc, boxTyPair, box, rewriter, kAttrAllocatable);472 rewriter.replaceOp(boxisalloc, check);473 return mlir::success();474 }475};476 477/// Lower `fir.box_isarray` to a sequence of operations to determine if the478/// boxed is an array.479struct BoxIsArrayOpConversion : public fir::FIROpConversion<fir::BoxIsArrayOp> {480 using FIROpConversion::FIROpConversion;481 482 llvm::LogicalResult483 matchAndRewrite(fir::BoxIsArrayOp boxisarray, OpAdaptor adaptor,484 mlir::ConversionPatternRewriter &rewriter) const override {485 mlir::Value a = adaptor.getOperands()[0];486 auto loc = boxisarray.getLoc();487 TypePair boxTyPair = getBoxTypePair(boxisarray.getVal().getType());488 mlir::Value rank = getRankFromBox(loc, boxTyPair, a, rewriter);489 mlir::Value c0 = fir::genConstantIndex(loc, rank.getType(), rewriter, 0);490 rewriter.replaceOpWithNewOp<mlir::LLVM::ICmpOp>(491 boxisarray, mlir::LLVM::ICmpPredicate::ne, rank, c0);492 return mlir::success();493 }494};495 496/// Lower `fir.box_isptr` to a sequence of operations to determined if the497/// boxed value was from a POINTER entity.498struct BoxIsPtrOpConversion : public fir::FIROpConversion<fir::BoxIsPtrOp> {499 using FIROpConversion::FIROpConversion;500 501 llvm::LogicalResult502 matchAndRewrite(fir::BoxIsPtrOp boxisptr, OpAdaptor adaptor,503 mlir::ConversionPatternRewriter &rewriter) const override {504 mlir::Value box = adaptor.getOperands()[0];505 auto loc = boxisptr.getLoc();506 TypePair boxTyPair = getBoxTypePair(boxisptr.getVal().getType());507 mlir::Value check =508 genBoxAttributeCheck(loc, boxTyPair, box, rewriter, kAttrPointer);509 rewriter.replaceOp(boxisptr, check);510 return mlir::success();511 }512};513 514/// Lower `fir.box_rank` to the sequence of operation to extract the rank from515/// the box.516struct BoxRankOpConversion : public fir::FIROpConversion<fir::BoxRankOp> {517 using FIROpConversion::FIROpConversion;518 519 llvm::LogicalResult520 matchAndRewrite(fir::BoxRankOp boxrank, OpAdaptor adaptor,521 mlir::ConversionPatternRewriter &rewriter) const override {522 mlir::Value a = adaptor.getOperands()[0];523 auto loc = boxrank.getLoc();524 mlir::Type ty = convertType(boxrank.getType());525 TypePair boxTyPair =526 getBoxTypePair(fir::unwrapRefType(boxrank.getBox().getType()));527 mlir::Value rank = getRankFromBox(loc, boxTyPair, a, rewriter);528 mlir::Value result = integerCast(loc, rewriter, ty, rank);529 rewriter.replaceOp(boxrank, result);530 return mlir::success();531 }532};533 534/// Lower `fir.boxproc_host` operation. Extracts the host pointer from the535/// boxproc.536/// TODO: Part of supporting Fortran 2003 procedure pointers.537struct BoxProcHostOpConversion538 : public fir::FIROpConversion<fir::BoxProcHostOp> {539 using FIROpConversion::FIROpConversion;540 541 llvm::LogicalResult542 matchAndRewrite(fir::BoxProcHostOp boxprochost, OpAdaptor adaptor,543 mlir::ConversionPatternRewriter &rewriter) const override {544 TODO(boxprochost.getLoc(), "fir.boxproc_host codegen");545 return mlir::failure();546 }547};548 549/// Lower `fir.box_tdesc` to the sequence of operations to extract the type550/// descriptor from the box.551struct BoxTypeDescOpConversion552 : public fir::FIROpConversion<fir::BoxTypeDescOp> {553 using FIROpConversion::FIROpConversion;554 555 llvm::LogicalResult556 matchAndRewrite(fir::BoxTypeDescOp boxtypedesc, OpAdaptor adaptor,557 mlir::ConversionPatternRewriter &rewriter) const override {558 mlir::Value box = adaptor.getOperands()[0];559 TypePair boxTyPair = getBoxTypePair(boxtypedesc.getBox().getType());560 auto typeDescAddr =561 loadTypeDescAddress(boxtypedesc.getLoc(), boxTyPair, box, rewriter);562 rewriter.replaceOp(boxtypedesc, typeDescAddr);563 return mlir::success();564 }565};566 567/// Lower `fir.box_typecode` to a sequence of operations to extract the type568/// code in the boxed value.569struct BoxTypeCodeOpConversion570 : public fir::FIROpConversion<fir::BoxTypeCodeOp> {571 using FIROpConversion::FIROpConversion;572 573 llvm::LogicalResult574 matchAndRewrite(fir::BoxTypeCodeOp op, OpAdaptor adaptor,575 mlir::ConversionPatternRewriter &rewriter) const override {576 mlir::Value box = adaptor.getOperands()[0];577 auto loc = box.getLoc();578 auto ty = convertType(op.getType());579 TypePair boxTyPair = getBoxTypePair(op.getBox().getType());580 auto typeCode =581 getValueFromBox(loc, boxTyPair, box, ty, rewriter, kTypePosInBox);582 rewriter.replaceOp(op, typeCode);583 return mlir::success();584 }585};586 587/// Lower `fir.string_lit` to LLVM IR dialect operation.588struct StringLitOpConversion : public fir::FIROpConversion<fir::StringLitOp> {589 using FIROpConversion::FIROpConversion;590 591 llvm::LogicalResult592 matchAndRewrite(fir::StringLitOp constop, OpAdaptor adaptor,593 mlir::ConversionPatternRewriter &rewriter) const override {594 auto ty = convertType(constop.getType());595 auto attr = constop.getValue();596 if (mlir::isa<mlir::StringAttr>(attr)) {597 rewriter.replaceOpWithNewOp<mlir::LLVM::ConstantOp>(constop, ty, attr);598 return mlir::success();599 }600 601 auto charTy = mlir::cast<fir::CharacterType>(constop.getType());602 unsigned bits = lowerTy().characterBitsize(charTy);603 mlir::Type intTy = rewriter.getIntegerType(bits);604 mlir::Location loc = constop.getLoc();605 mlir::Value cst = mlir::LLVM::UndefOp::create(rewriter, loc, ty);606 if (auto arr = mlir::dyn_cast<mlir::DenseElementsAttr>(attr)) {607 cst = mlir::LLVM::ConstantOp::create(rewriter, loc, ty, arr);608 } else if (auto arr = mlir::dyn_cast<mlir::ArrayAttr>(attr)) {609 for (auto a : llvm::enumerate(arr.getValue())) {610 // convert each character to a precise bitsize611 auto elemAttr = mlir::IntegerAttr::get(612 intTy,613 mlir::cast<mlir::IntegerAttr>(a.value()).getValue().zextOrTrunc(614 bits));615 auto elemCst =616 mlir::LLVM::ConstantOp::create(rewriter, loc, intTy, elemAttr);617 cst = mlir::LLVM::InsertValueOp::create(rewriter, loc, cst, elemCst,618 a.index());619 }620 } else {621 return mlir::failure();622 }623 rewriter.replaceOp(constop, cst);624 return mlir::success();625 }626};627 628/// `fir.call` -> `llvm.call`629struct CallOpConversion : public fir::FIROpConversion<fir::CallOp> {630 using FIROpConversion::FIROpConversion;631 632 llvm::LogicalResult633 matchAndRewrite(fir::CallOp call, OpAdaptor adaptor,634 mlir::ConversionPatternRewriter &rewriter) const override {635 llvm::SmallVector<mlir::Type> resultTys;636 mlir::Attribute memAttr =637 call->getAttr(fir::FIROpsDialect::getFirCallMemoryAttrName());638 if (memAttr)639 call->removeAttr(fir::FIROpsDialect::getFirCallMemoryAttrName());640 641 for (auto r : call.getResults())642 resultTys.push_back(convertType(r.getType()));643 // Convert arith::FastMathFlagsAttr to LLVM::FastMathFlagsAttr.644 mlir::arith::AttrConvertFastMathToLLVM<fir::CallOp, mlir::LLVM::CallOp>645 attrConvert(call);646 auto llvmCall = rewriter.replaceOpWithNewOp<mlir::LLVM::CallOp>(647 call, resultTys, adaptor.getOperands(),648 addLLVMOpBundleAttrs(rewriter, attrConvert.getAttrs(),649 adaptor.getOperands().size()));650 if (mlir::ArrayAttr argAttrsArray = call.getArgAttrsAttr()) {651 // sret and byval type needs to be converted.652 auto convertTypeAttr = [&](const mlir::NamedAttribute &attr) {653 return mlir::TypeAttr::get(convertType(654 llvm::cast<mlir::TypeAttr>(attr.getValue()).getValue()));655 };656 llvm::SmallVector<mlir::Attribute> newArgAttrsArray;657 for (auto argAttrs : argAttrsArray) {658 llvm::SmallVector<mlir::NamedAttribute> convertedAttrs;659 for (const mlir::NamedAttribute &attr :660 llvm::cast<mlir::DictionaryAttr>(argAttrs)) {661 if (attr.getName().getValue() ==662 mlir::LLVM::LLVMDialect::getByValAttrName()) {663 convertedAttrs.push_back(rewriter.getNamedAttr(664 mlir::LLVM::LLVMDialect::getByValAttrName(),665 convertTypeAttr(attr)));666 } else if (attr.getName().getValue() ==667 mlir::LLVM::LLVMDialect::getStructRetAttrName()) {668 convertedAttrs.push_back(rewriter.getNamedAttr(669 mlir::LLVM::LLVMDialect::getStructRetAttrName(),670 convertTypeAttr(attr)));671 } else {672 convertedAttrs.push_back(attr);673 }674 }675 newArgAttrsArray.emplace_back(676 mlir::DictionaryAttr::get(rewriter.getContext(), convertedAttrs));677 }678 llvmCall.setArgAttrsAttr(rewriter.getArrayAttr(newArgAttrsArray));679 }680 if (mlir::ArrayAttr resAttrs = call.getResAttrsAttr())681 llvmCall.setResAttrsAttr(resAttrs);682 683 if (auto inlineAttr = call.getInlineAttrAttr()) {684 llvmCall->removeAttr("inline_attr");685 if (inlineAttr.getValue() == fir::FortranInlineEnum::no_inline) {686 llvmCall.setNoInlineAttr(rewriter.getUnitAttr());687 } else if (inlineAttr.getValue() == fir::FortranInlineEnum::inline_hint) {688 llvmCall.setInlineHintAttr(rewriter.getUnitAttr());689 } else if (inlineAttr.getValue() ==690 fir::FortranInlineEnum::always_inline) {691 llvmCall.setAlwaysInlineAttr(rewriter.getUnitAttr());692 }693 }694 695 if (std::optional<mlir::ArrayAttr> optionalAccessGroups =696 call.getAccessGroups())697 llvmCall.setAccessGroups(*optionalAccessGroups);698 699 if (memAttr)700 llvmCall.setMemoryEffectsAttr(701 mlir::cast<mlir::LLVM::MemoryEffectsAttr>(memAttr));702 return mlir::success();703 }704};705} // namespace706 707static mlir::Type getComplexEleTy(mlir::Type complex) {708 return mlir::cast<mlir::ComplexType>(complex).getElementType();709}710 711namespace {712/// Compare complex values713///714/// Per 10.1, the only comparisons available are .EQ. (oeq) and .NE. (une).715///716/// For completeness, all other comparison are done on the real component only.717struct CmpcOpConversion : public fir::FIROpConversion<fir::CmpcOp> {718 using FIROpConversion::FIROpConversion;719 720 llvm::LogicalResult721 matchAndRewrite(fir::CmpcOp cmp, OpAdaptor adaptor,722 mlir::ConversionPatternRewriter &rewriter) const override {723 mlir::ValueRange operands = adaptor.getOperands();724 mlir::Type resTy = convertType(cmp.getType());725 mlir::Location loc = cmp.getLoc();726 mlir::LLVM::FastmathFlags fmf =727 mlir::arith::convertArithFastMathFlagsToLLVM(cmp.getFastmath());728 mlir::LLVM::FCmpPredicate pred =729 static_cast<mlir::LLVM::FCmpPredicate>(cmp.getPredicate());730 auto rcp = mlir::LLVM::FCmpOp::create(731 rewriter, loc, resTy, pred,732 mlir::LLVM::ExtractValueOp::create(rewriter, loc, operands[0], 0),733 mlir::LLVM::ExtractValueOp::create(rewriter, loc, operands[1], 0), fmf);734 auto icp = mlir::LLVM::FCmpOp::create(735 rewriter, loc, resTy, pred,736 mlir::LLVM::ExtractValueOp::create(rewriter, loc, operands[0], 1),737 mlir::LLVM::ExtractValueOp::create(rewriter, loc, operands[1], 1), fmf);738 llvm::SmallVector<mlir::Value, 2> cp = {rcp, icp};739 switch (cmp.getPredicate()) {740 case mlir::arith::CmpFPredicate::OEQ: // .EQ.741 rewriter.replaceOpWithNewOp<mlir::LLVM::AndOp>(cmp, resTy, cp);742 break;743 case mlir::arith::CmpFPredicate::UNE: // .NE.744 rewriter.replaceOpWithNewOp<mlir::LLVM::OrOp>(cmp, resTy, cp);745 break;746 default:747 rewriter.replaceOp(cmp, rcp.getResult());748 break;749 }750 return mlir::success();751 }752};753 754/// fir.volatile_cast is only useful at the fir level. Once we lower to LLVM,755/// volatility is described by setting volatile attributes on the LLVM ops.756struct VolatileCastOpConversion757 : public fir::FIROpConversion<fir::VolatileCastOp> {758 using FIROpConversion::FIROpConversion;759 760 llvm::LogicalResult761 matchAndRewrite(fir::VolatileCastOp volatileCast, OpAdaptor adaptor,762 mlir::ConversionPatternRewriter &rewriter) const override {763 rewriter.replaceOp(volatileCast, adaptor.getOperands()[0]);764 return mlir::success();765 }766};767 768/// Lower `fir.assumed_size_extent` to constant -1 of index type.769struct AssumedSizeExtentOpConversion770 : public fir::FIROpConversion<fir::AssumedSizeExtentOp> {771 using FIROpConversion::FIROpConversion;772 773 llvm::LogicalResult774 matchAndRewrite(fir::AssumedSizeExtentOp op, OpAdaptor,775 mlir::ConversionPatternRewriter &rewriter) const override {776 mlir::Location loc = op.getLoc();777 mlir::Type ity = lowerTy().indexType();778 auto cst = fir::genConstantIndex(loc, ity, rewriter, -1);779 rewriter.replaceOp(op, cst.getResult());780 return mlir::success();781 }782};783 784/// Lower `fir.is_assumed_size_extent` to integer equality with -1.785struct IsAssumedSizeExtentOpConversion786 : public fir::FIROpConversion<fir::IsAssumedSizeExtentOp> {787 using FIROpConversion::FIROpConversion;788 789 llvm::LogicalResult790 matchAndRewrite(fir::IsAssumedSizeExtentOp op, OpAdaptor adaptor,791 mlir::ConversionPatternRewriter &rewriter) const override {792 mlir::Location loc = op.getLoc();793 mlir::Value val = adaptor.getVal();794 mlir::Type valTy = val.getType();795 // Create constant -1 of the operand type.796 auto negOneAttr = rewriter.getIntegerAttr(valTy, -1);797 auto negOne =798 mlir::LLVM::ConstantOp::create(rewriter, loc, valTy, negOneAttr);799 auto cmp = mlir::LLVM::ICmpOp::create(800 rewriter, loc, mlir::LLVM::ICmpPredicate::eq, val, negOne);801 rewriter.replaceOp(op, cmp.getResult());802 return mlir::success();803 }804};805 806/// convert value of from-type to value of to-type807struct ConvertOpConversion : public fir::FIROpConversion<fir::ConvertOp> {808 using FIROpConversion::FIROpConversion;809 810 static bool isFloatingPointTy(mlir::Type ty) {811 return mlir::isa<mlir::FloatType>(ty);812 }813 814 llvm::LogicalResult815 matchAndRewrite(fir::ConvertOp convert, OpAdaptor adaptor,816 mlir::ConversionPatternRewriter &rewriter) const override {817 auto fromFirTy = convert.getValue().getType();818 auto toFirTy = convert.getRes().getType();819 auto fromTy = convertType(fromFirTy);820 auto toTy = convertType(toFirTy);821 mlir::Value op0 = adaptor.getOperands()[0];822 823 if (fromFirTy == toFirTy) {824 rewriter.replaceOp(convert, op0);825 return mlir::success();826 }827 828 auto loc = convert.getLoc();829 auto i1Type = mlir::IntegerType::get(convert.getContext(), 1);830 831 if (mlir::isa<fir::RecordType>(toFirTy)) {832 // Convert to compatible BIND(C) record type.833 // Double check that the record types are compatible (it should have834 // already been checked by the verifier).835 assert(mlir::cast<fir::RecordType>(fromFirTy).getTypeList() ==836 mlir::cast<fir::RecordType>(toFirTy).getTypeList() &&837 "incompatible record types");838 839 auto toStTy = mlir::cast<mlir::LLVM::LLVMStructType>(toTy);840 mlir::Value val = mlir::LLVM::UndefOp::create(rewriter, loc, toStTy);841 auto indexTypeMap = toStTy.getSubelementIndexMap();842 assert(indexTypeMap.has_value() && "invalid record type");843 844 for (auto [attr, type] : indexTypeMap.value()) {845 int64_t index = mlir::cast<mlir::IntegerAttr>(attr).getInt();846 auto extVal =847 mlir::LLVM::ExtractValueOp::create(rewriter, loc, op0, index);848 val = mlir::LLVM::InsertValueOp::create(rewriter, loc, val, extVal,849 index);850 }851 852 rewriter.replaceOp(convert, val);853 return mlir::success();854 }855 856 if (mlir::isa<fir::LogicalType>(fromFirTy) ||857 mlir::isa<fir::LogicalType>(toFirTy)) {858 // By specification fir::LogicalType value may be any number,859 // where non-zero value represents .true. and zero value represents860 // .false.861 //862 // integer<->logical conversion requires value normalization.863 // Conversion from wide logical to narrow logical must set the result864 // to non-zero iff the input is non-zero - the easiest way to implement865 // it is to compare the input agains zero and set the result to866 // the canonical 0/1.867 // Conversion from narrow logical to wide logical may be implemented868 // as a zero or sign extension of the input, but it may use value869 // normalization as well.870 if (!mlir::isa<mlir::IntegerType>(fromTy) ||871 !mlir::isa<mlir::IntegerType>(toTy))872 return mlir::emitError(loc)873 << "unsupported types for logical conversion: " << fromTy874 << " -> " << toTy;875 876 // Do folding for constant inputs.877 if (auto constVal = fir::getIntIfConstant(op0)) {878 mlir::Value normVal =879 fir::genConstantIndex(loc, toTy, rewriter, *constVal ? 1 : 0);880 rewriter.replaceOp(convert, normVal);881 return mlir::success();882 }883 884 // If the input is i1, then we can just zero extend it, and885 // the result will be normalized.886 if (fromTy == i1Type) {887 rewriter.replaceOpWithNewOp<mlir::LLVM::ZExtOp>(convert, toTy, op0);888 return mlir::success();889 }890 891 // Compare the input with zero.892 mlir::Value zero = fir::genConstantIndex(loc, fromTy, rewriter, 0);893 auto isTrue = mlir::LLVM::ICmpOp::create(894 rewriter, loc, mlir::LLVM::ICmpPredicate::ne, op0, zero);895 896 // Zero extend the i1 isTrue result to the required type (unless it is i1897 // itself).898 if (toTy != i1Type)899 rewriter.replaceOpWithNewOp<mlir::LLVM::ZExtOp>(convert, toTy, isTrue);900 else901 rewriter.replaceOp(convert, isTrue.getResult());902 903 return mlir::success();904 }905 906 if (fromTy == toTy) {907 rewriter.replaceOp(convert, op0);908 return mlir::success();909 }910 auto convertFpToFp = [&](mlir::Value val, unsigned fromBits,911 unsigned toBits, mlir::Type toTy) -> mlir::Value {912 if (fromBits == toBits) {913 // TODO: Converting between two floating-point representations with the914 // same bitwidth is not allowed for now.915 mlir::emitError(loc,916 "cannot implicitly convert between two floating-point "917 "representations of the same bitwidth");918 return {};919 }920 if (fromBits > toBits)921 return mlir::LLVM::FPTruncOp::create(rewriter, loc, toTy, val);922 return mlir::LLVM::FPExtOp::create(rewriter, loc, toTy, val);923 };924 // Complex to complex conversion.925 if (fir::isa_complex(fromFirTy) && fir::isa_complex(toFirTy)) {926 // Special case: handle the conversion of a complex such that both the927 // real and imaginary parts are converted together.928 auto ty = convertType(getComplexEleTy(convert.getValue().getType()));929 auto rp = mlir::LLVM::ExtractValueOp::create(rewriter, loc, op0, 0);930 auto ip = mlir::LLVM::ExtractValueOp::create(rewriter, loc, op0, 1);931 auto nt = convertType(getComplexEleTy(convert.getRes().getType()));932 auto fromBits = mlir::LLVM::getPrimitiveTypeSizeInBits(ty);933 auto toBits = mlir::LLVM::getPrimitiveTypeSizeInBits(nt);934 auto rc = convertFpToFp(rp, fromBits, toBits, nt);935 auto ic = convertFpToFp(ip, fromBits, toBits, nt);936 auto un = mlir::LLVM::UndefOp::create(rewriter, loc, toTy);937 llvm::SmallVector<int64_t> pos{0};938 auto i1 = mlir::LLVM::InsertValueOp::create(rewriter, loc, un, rc, pos);939 rewriter.replaceOpWithNewOp<mlir::LLVM::InsertValueOp>(convert, i1, ic,940 1);941 return mlir::success();942 }943 944 // Floating point to floating point conversion.945 if (isFloatingPointTy(fromTy)) {946 if (isFloatingPointTy(toTy)) {947 auto fromBits = mlir::LLVM::getPrimitiveTypeSizeInBits(fromTy);948 auto toBits = mlir::LLVM::getPrimitiveTypeSizeInBits(toTy);949 auto v = convertFpToFp(op0, fromBits, toBits, toTy);950 rewriter.replaceOp(convert, v);951 return mlir::success();952 }953 if (mlir::isa<mlir::IntegerType>(toTy)) {954 // NOTE: We are checking the fir type here because toTy is an LLVM type955 // which is signless, and we need to use the intrinsic that matches the956 // sign of the output in fir.957 if (toFirTy.isUnsignedInteger()) {958 auto intrinsicName =959 mlir::StringAttr::get(convert.getContext(), "llvm.fptoui.sat");960 rewriter.replaceOpWithNewOp<mlir::LLVM::CallIntrinsicOp>(961 convert, toTy, intrinsicName, op0);962 } else {963 auto intrinsicName =964 mlir::StringAttr::get(convert.getContext(), "llvm.fptosi.sat");965 rewriter.replaceOpWithNewOp<mlir::LLVM::CallIntrinsicOp>(966 convert, toTy, intrinsicName, op0);967 }968 return mlir::success();969 }970 } else if (mlir::isa<mlir::IntegerType>(fromTy)) {971 // Integer to integer conversion.972 if (mlir::isa<mlir::IntegerType>(toTy)) {973 auto fromBits = mlir::LLVM::getPrimitiveTypeSizeInBits(fromTy);974 auto toBits = mlir::LLVM::getPrimitiveTypeSizeInBits(toTy);975 assert(fromBits != toBits);976 if (fromBits > toBits) {977 rewriter.replaceOpWithNewOp<mlir::LLVM::TruncOp>(convert, toTy, op0);978 return mlir::success();979 }980 if (fromFirTy == i1Type || fromFirTy.isUnsignedInteger()) {981 rewriter.replaceOpWithNewOp<mlir::LLVM::ZExtOp>(convert, toTy, op0);982 return mlir::success();983 }984 rewriter.replaceOpWithNewOp<mlir::LLVM::SExtOp>(convert, toTy, op0);985 return mlir::success();986 }987 // Integer to floating point conversion.988 if (isFloatingPointTy(toTy)) {989 if (fromTy.isUnsignedInteger())990 rewriter.replaceOpWithNewOp<mlir::LLVM::UIToFPOp>(convert, toTy, op0);991 else992 rewriter.replaceOpWithNewOp<mlir::LLVM::SIToFPOp>(convert, toTy, op0);993 return mlir::success();994 }995 // Integer to pointer conversion.996 if (mlir::isa<mlir::LLVM::LLVMPointerType>(toTy)) {997 rewriter.replaceOpWithNewOp<mlir::LLVM::IntToPtrOp>(convert, toTy, op0);998 return mlir::success();999 }1000 } else if (mlir::isa<mlir::LLVM::LLVMPointerType>(fromTy)) {1001 // Pointer to integer conversion.1002 if (mlir::isa<mlir::IntegerType>(toTy)) {1003 rewriter.replaceOpWithNewOp<mlir::LLVM::PtrToIntOp>(convert, toTy, op0);1004 return mlir::success();1005 }1006 // Pointer to pointer conversion.1007 if (mlir::isa<mlir::LLVM::LLVMPointerType>(toTy)) {1008 rewriter.replaceOpWithNewOp<mlir::LLVM::BitcastOp>(convert, toTy, op0);1009 return mlir::success();1010 }1011 }1012 return emitError(loc) << "cannot convert " << fromTy << " to " << toTy;1013 }1014};1015 1016/// `fir.type_info` operation has no specific CodeGen. The operation is1017/// only used to carry information during FIR to FIR passes. It may be used1018/// in the future to generate the runtime type info data structures instead1019/// of generating them in lowering.1020struct TypeInfoOpConversion : public fir::FIROpConversion<fir::TypeInfoOp> {1021 using FIROpConversion::FIROpConversion;1022 1023 llvm::LogicalResult1024 matchAndRewrite(fir::TypeInfoOp op, OpAdaptor,1025 mlir::ConversionPatternRewriter &rewriter) const override {1026 rewriter.eraseOp(op);1027 return mlir::success();1028 }1029};1030 1031/// `fir.dt_entry` operation has no specific CodeGen. The operation is only used1032/// to carry information during FIR to FIR passes.1033struct DTEntryOpConversion : public fir::FIROpConversion<fir::DTEntryOp> {1034 using FIROpConversion::FIROpConversion;1035 1036 llvm::LogicalResult1037 matchAndRewrite(fir::DTEntryOp op, OpAdaptor,1038 mlir::ConversionPatternRewriter &rewriter) const override {1039 rewriter.eraseOp(op);1040 return mlir::success();1041 }1042};1043 1044/// Lower `fir.global_len` operation.1045struct GlobalLenOpConversion : public fir::FIROpConversion<fir::GlobalLenOp> {1046 using FIROpConversion::FIROpConversion;1047 1048 llvm::LogicalResult1049 matchAndRewrite(fir::GlobalLenOp globalLen, OpAdaptor adaptor,1050 mlir::ConversionPatternRewriter &rewriter) const override {1051 TODO(globalLen.getLoc(), "fir.global_len codegen");1052 return mlir::failure();1053 }1054};1055 1056/// Lower fir.len_param_index1057struct LenParamIndexOpConversion1058 : public fir::FIROpConversion<fir::LenParamIndexOp> {1059 using FIROpConversion::FIROpConversion;1060 1061 // FIXME: this should be specialized by the runtime target1062 llvm::LogicalResult1063 matchAndRewrite(fir::LenParamIndexOp lenp, OpAdaptor,1064 mlir::ConversionPatternRewriter &rewriter) const override {1065 TODO(lenp.getLoc(), "fir.len_param_index codegen");1066 }1067};1068 1069/// Convert `!fir.emboxchar<!fir.char<KIND, ?>, #n>` into a sequence of1070/// instructions that generate `!llvm.struct<(ptr<ik>, i64)>`. The 1st element1071/// in this struct is a pointer. Its type is determined from `KIND`. The 2nd1072/// element is the length of the character buffer (`#n`).1073struct EmboxCharOpConversion : public fir::FIROpConversion<fir::EmboxCharOp> {1074 using FIROpConversion::FIROpConversion;1075 1076 llvm::LogicalResult1077 matchAndRewrite(fir::EmboxCharOp emboxChar, OpAdaptor adaptor,1078 mlir::ConversionPatternRewriter &rewriter) const override {1079 mlir::ValueRange operands = adaptor.getOperands();1080 1081 mlir::Value charBuffer = operands[0];1082 mlir::Value charBufferLen = operands[1];1083 1084 mlir::Location loc = emboxChar.getLoc();1085 mlir::Type llvmStructTy = convertType(emboxChar.getType());1086 auto llvmStruct = mlir::LLVM::UndefOp::create(rewriter, loc, llvmStructTy);1087 1088 mlir::Type lenTy =1089 mlir::cast<mlir::LLVM::LLVMStructType>(llvmStructTy).getBody()[1];1090 mlir::Value lenAfterCast = integerCast(loc, rewriter, lenTy, charBufferLen);1091 1092 mlir::Type addrTy =1093 mlir::cast<mlir::LLVM::LLVMStructType>(llvmStructTy).getBody()[0];1094 if (addrTy != charBuffer.getType())1095 charBuffer =1096 mlir::LLVM::BitcastOp::create(rewriter, loc, addrTy, charBuffer);1097 1098 llvm::SmallVector<int64_t> pos{0};1099 auto insertBufferOp = mlir::LLVM::InsertValueOp::create(1100 rewriter, loc, llvmStruct, charBuffer, pos);1101 rewriter.replaceOpWithNewOp<mlir::LLVM::InsertValueOp>(1102 emboxChar, insertBufferOp, lenAfterCast, 1);1103 1104 return mlir::success();1105 }1106};1107} // namespace1108 1109template <typename ModuleOp>1110static mlir::SymbolRefAttr1111getMallocInModule(ModuleOp mod, fir::AllocMemOp op,1112 mlir::ConversionPatternRewriter &rewriter,1113 mlir::Type indexType) {1114 static constexpr char mallocName[] = "malloc";1115 if (auto mallocFunc =1116 mod.template lookupSymbol<mlir::LLVM::LLVMFuncOp>(mallocName))1117 return mlir::SymbolRefAttr::get(mallocFunc);1118 if (auto userMalloc =1119 mod.template lookupSymbol<mlir::func::FuncOp>(mallocName))1120 return mlir::SymbolRefAttr::get(userMalloc);1121 1122 mlir::OpBuilder moduleBuilder(mod.getBodyRegion());1123 auto mallocDecl = mlir::LLVM::LLVMFuncOp::create(1124 moduleBuilder, op.getLoc(), mallocName,1125 mlir::LLVM::LLVMFunctionType::get(getLlvmPtrType(op.getContext()),1126 indexType,1127 /*isVarArg=*/false));1128 return mlir::SymbolRefAttr::get(mallocDecl);1129}1130 1131/// Return the LLVMFuncOp corresponding to the standard malloc call.1132static mlir::SymbolRefAttr getMalloc(fir::AllocMemOp op,1133 mlir::ConversionPatternRewriter &rewriter,1134 mlir::Type indexType) {1135 if (auto mod = op->getParentOfType<mlir::gpu::GPUModuleOp>())1136 return getMallocInModule(mod, op, rewriter, indexType);1137 auto mod = op->getParentOfType<mlir::ModuleOp>();1138 return getMallocInModule(mod, op, rewriter, indexType);1139}1140 1141/// Return value of the stride in bytes between adjacent elements1142/// of LLVM type \p llTy. The result is returned as a value of1143/// \p idxTy integer type.1144static mlir::Value1145genTypeStrideInBytes(mlir::Location loc, mlir::Type idxTy,1146 mlir::ConversionPatternRewriter &rewriter, mlir::Type llTy,1147 const mlir::DataLayout &dataLayout) {1148 // Create a pointer type and use computeElementDistance().1149 return fir::computeElementDistance(loc, llTy, idxTy, rewriter, dataLayout);1150}1151 1152namespace {1153/// Lower a `fir.allocmem` instruction into `llvm.call @malloc`1154struct AllocMemOpConversion : public fir::FIROpConversion<fir::AllocMemOp> {1155 using FIROpConversion::FIROpConversion;1156 1157 llvm::LogicalResult1158 matchAndRewrite(fir::AllocMemOp heap, OpAdaptor adaptor,1159 mlir::ConversionPatternRewriter &rewriter) const override {1160 mlir::Type heapTy = heap.getType();1161 mlir::Location loc = heap.getLoc();1162 auto ity = lowerTy().indexType();1163 mlir::Type dataTy = fir::unwrapRefType(heapTy);1164 mlir::Type llvmObjectTy = convertObjectType(dataTy);1165 if (fir::isRecordWithTypeParameters(fir::unwrapSequenceType(dataTy)))1166 TODO(loc, "fir.allocmem codegen of derived type with length parameters");1167 mlir::Value size = genTypeSizeInBytes(loc, ity, rewriter, llvmObjectTy);1168 if (auto scaleSize =1169 fir::genAllocationScaleSize(loc, heap.getInType(), ity, rewriter))1170 size = mlir::LLVM::MulOp::create(rewriter, loc, ity, size, scaleSize);1171 for (mlir::Value opnd : adaptor.getOperands())1172 size = mlir::LLVM::MulOp::create(rewriter, loc, ity, size,1173 integerCast(loc, rewriter, ity, opnd));1174 1175 // As the return value of malloc(0) is implementation defined, allocate one1176 // byte to ensure the allocation status being true. This behavior aligns to1177 // what the runtime has.1178 mlir::Value zero = fir::genConstantIndex(loc, ity, rewriter, 0);1179 mlir::Value one = fir::genConstantIndex(loc, ity, rewriter, 1);1180 mlir::Value cmp = mlir::LLVM::ICmpOp::create(1181 rewriter, loc, mlir::LLVM::ICmpPredicate::sgt, size, zero);1182 size = mlir::LLVM::SelectOp::create(rewriter, loc, cmp, size, one);1183 1184 auto mallocTyWidth = lowerTy().getIndexTypeBitwidth();1185 auto mallocTy =1186 mlir::IntegerType::get(rewriter.getContext(), mallocTyWidth);1187 if (mallocTyWidth != ity.getIntOrFloatBitWidth())1188 size = integerCast(loc, rewriter, mallocTy, size);1189 heap->setAttr("callee", getMalloc(heap, rewriter, mallocTy));1190 rewriter.replaceOpWithNewOp<mlir::LLVM::CallOp>(1191 heap, ::getLlvmPtrType(heap.getContext()), size,1192 addLLVMOpBundleAttrs(rewriter, heap->getAttrs(), 1));1193 return mlir::success();1194 }1195 1196 /// Compute the allocation size in bytes of the element type of1197 /// \p llTy pointer type. The result is returned as a value of \p idxTy1198 /// integer type.1199 mlir::Value genTypeSizeInBytes(mlir::Location loc, mlir::Type idxTy,1200 mlir::ConversionPatternRewriter &rewriter,1201 mlir::Type llTy) const {1202 return fir::computeElementDistance(loc, llTy, idxTy, rewriter,1203 getDataLayout());1204 }1205};1206} // namespace1207 1208/// Return the LLVMFuncOp corresponding to the standard free call.1209template <typename ModuleOp>1210static mlir::SymbolRefAttr1211getFreeInModule(ModuleOp mod, fir::FreeMemOp op,1212 mlir::ConversionPatternRewriter &rewriter) {1213 static constexpr char freeName[] = "free";1214 // Check if free already defined in the module.1215 if (auto freeFunc =1216 mod.template lookupSymbol<mlir::LLVM::LLVMFuncOp>(freeName))1217 return mlir::SymbolRefAttr::get(freeFunc);1218 if (auto freeDefinedByUser =1219 mod.template lookupSymbol<mlir::func::FuncOp>(freeName))1220 return mlir::SymbolRefAttr::get(freeDefinedByUser);1221 // Create llvm declaration for free.1222 mlir::OpBuilder moduleBuilder(mod.getBodyRegion());1223 auto voidType = mlir::LLVM::LLVMVoidType::get(op.getContext());1224 auto freeDecl = mlir::LLVM::LLVMFuncOp::create(1225 moduleBuilder, rewriter.getUnknownLoc(), freeName,1226 mlir::LLVM::LLVMFunctionType::get(voidType,1227 getLlvmPtrType(op.getContext()),1228 /*isVarArg=*/false));1229 return mlir::SymbolRefAttr::get(freeDecl);1230}1231 1232static mlir::SymbolRefAttr getFree(fir::FreeMemOp op,1233 mlir::ConversionPatternRewriter &rewriter) {1234 if (auto mod = op->getParentOfType<mlir::gpu::GPUModuleOp>())1235 return getFreeInModule(mod, op, rewriter);1236 auto mod = op->getParentOfType<mlir::ModuleOp>();1237 return getFreeInModule(mod, op, rewriter);1238}1239 1240static unsigned getDimension(mlir::LLVM::LLVMArrayType ty) {1241 unsigned result = 1;1242 for (auto eleTy =1243 mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(ty.getElementType());1244 eleTy; eleTy = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(1245 eleTy.getElementType()))1246 ++result;1247 return result;1248}1249 1250namespace {1251/// Lower a `fir.freemem` instruction into `llvm.call @free`1252struct FreeMemOpConversion : public fir::FIROpConversion<fir::FreeMemOp> {1253 using FIROpConversion::FIROpConversion;1254 1255 llvm::LogicalResult1256 matchAndRewrite(fir::FreeMemOp freemem, OpAdaptor adaptor,1257 mlir::ConversionPatternRewriter &rewriter) const override {1258 mlir::Location loc = freemem.getLoc();1259 freemem->setAttr("callee", getFree(freemem, rewriter));1260 mlir::LLVM::CallOp::create(1261 rewriter, loc, mlir::TypeRange{},1262 mlir::ValueRange{adaptor.getHeapref()},1263 addLLVMOpBundleAttrs(rewriter, freemem->getAttrs(), 1));1264 rewriter.eraseOp(freemem);1265 return mlir::success();1266 }1267};1268} // namespace1269 1270// Convert subcomponent array indices from column-major to row-major ordering.1271static llvm::SmallVector<mlir::Value>1272convertSubcomponentIndices(mlir::Location loc, mlir::Type eleTy,1273 mlir::ValueRange indices,1274 mlir::Type *retTy = nullptr) {1275 llvm::SmallVector<mlir::Value> result;1276 llvm::SmallVector<mlir::Value> arrayIndices;1277 1278 auto appendArrayIndices = [&] {1279 if (arrayIndices.empty())1280 return;1281 std::reverse(arrayIndices.begin(), arrayIndices.end());1282 result.append(arrayIndices.begin(), arrayIndices.end());1283 arrayIndices.clear();1284 };1285 1286 for (mlir::Value index : indices) {1287 // Component indices can be field index to select a component, or array1288 // index, to select an element in an array component.1289 if (auto structTy = mlir::dyn_cast<mlir::LLVM::LLVMStructType>(eleTy)) {1290 std::int64_t cstIndex = getConstantIntValue(index);1291 assert(cstIndex < (int64_t)structTy.getBody().size() &&1292 "out-of-bounds struct field index");1293 eleTy = structTy.getBody()[cstIndex];1294 appendArrayIndices();1295 result.push_back(index);1296 } else if (auto arrayTy =1297 mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(eleTy)) {1298 eleTy = arrayTy.getElementType();1299 arrayIndices.push_back(index);1300 } else1301 fir::emitFatalError(loc, "Unexpected subcomponent type");1302 }1303 appendArrayIndices();1304 if (retTy)1305 *retTy = eleTy;1306 return result;1307}1308 1309static mlir::Value genSourceFile(mlir::Location loc, mlir::ModuleOp mod,1310 mlir::ConversionPatternRewriter &rewriter) {1311 auto ptrTy = mlir::LLVM::LLVMPointerType::get(rewriter.getContext());1312 if (auto flc = mlir::dyn_cast<mlir::FileLineColLoc>(loc)) {1313 auto fn = flc.getFilename().str() + '\0';1314 std::string globalName = fir::factory::uniqueCGIdent("cl", fn);1315 1316 if (auto g = mod.lookupSymbol<fir::GlobalOp>(globalName)) {1317 return mlir::LLVM::AddressOfOp::create(rewriter, loc, ptrTy, g.getName());1318 } else if (auto g = mod.lookupSymbol<mlir::LLVM::GlobalOp>(globalName)) {1319 return mlir::LLVM::AddressOfOp::create(rewriter, loc, ptrTy, g.getName());1320 }1321 1322 auto crtInsPt = rewriter.saveInsertionPoint();1323 rewriter.setInsertionPoint(mod.getBody(), mod.getBody()->end());1324 auto arrayTy = mlir::LLVM::LLVMArrayType::get(1325 mlir::IntegerType::get(rewriter.getContext(), 8), fn.size());1326 mlir::LLVM::GlobalOp globalOp = mlir::LLVM::GlobalOp::create(1327 rewriter, loc, arrayTy, /*constant=*/true,1328 mlir::LLVM::Linkage::Linkonce, globalName, mlir::Attribute());1329 1330 mlir::Region ®ion = globalOp.getInitializerRegion();1331 mlir::Block *block = rewriter.createBlock(®ion);1332 rewriter.setInsertionPoint(block, block->begin());1333 mlir::Value constValue = mlir::LLVM::ConstantOp::create(1334 rewriter, loc, arrayTy, rewriter.getStringAttr(fn));1335 mlir::LLVM::ReturnOp::create(rewriter, loc, constValue);1336 rewriter.restoreInsertionPoint(crtInsPt);1337 return mlir::LLVM::AddressOfOp::create(rewriter, loc, ptrTy,1338 globalOp.getName());1339 }1340 return mlir::LLVM::ZeroOp::create(rewriter, loc, ptrTy);1341}1342 1343static mlir::Value genSourceLine(mlir::Location loc,1344 mlir::ConversionPatternRewriter &rewriter) {1345 if (auto flc = mlir::dyn_cast<mlir::FileLineColLoc>(loc))1346 return mlir::LLVM::ConstantOp::create(rewriter, loc, rewriter.getI32Type(),1347 flc.getLine());1348 return mlir::LLVM::ConstantOp::create(rewriter, loc, rewriter.getI32Type(),1349 0);1350}1351 1352static mlir::Value1353genCUFAllocDescriptor(mlir::Location loc,1354 mlir::ConversionPatternRewriter &rewriter,1355 mlir::ModuleOp mod, fir::BaseBoxType boxTy,1356 const fir::LLVMTypeConverter &typeConverter) {1357 std::optional<mlir::DataLayout> dl =1358 fir::support::getOrSetMLIRDataLayout(mod, /*allowDefaultLayout=*/true);1359 if (!dl)1360 mlir::emitError(mod.getLoc(),1361 "module operation must carry a data layout attribute "1362 "to generate llvm IR from FIR");1363 1364 mlir::Value sourceFile = genSourceFile(loc, mod, rewriter);1365 mlir::Value sourceLine = genSourceLine(loc, rewriter);1366 1367 mlir::MLIRContext *ctx = mod.getContext();1368 1369 mlir::LLVM::LLVMPointerType llvmPointerType =1370 mlir::LLVM::LLVMPointerType::get(ctx);1371 mlir::Type llvmInt32Type = mlir::IntegerType::get(ctx, 32);1372 mlir::Type llvmIntPtrType =1373 mlir::IntegerType::get(ctx, typeConverter.getPointerBitwidth(0));1374 auto fctTy = mlir::LLVM::LLVMFunctionType::get(1375 llvmPointerType, {llvmIntPtrType, llvmPointerType, llvmInt32Type});1376 1377 auto llvmFunc = mod.lookupSymbol<mlir::LLVM::LLVMFuncOp>(1378 RTNAME_STRING(CUFAllocDescriptor));1379 auto funcFunc =1380 mod.lookupSymbol<mlir::func::FuncOp>(RTNAME_STRING(CUFAllocDescriptor));1381 if (!llvmFunc && !funcFunc) {1382 auto builder = mlir::OpBuilder::atBlockEnd(mod.getBody());1383 mlir::LLVM::LLVMFuncOp::create(builder, loc,1384 RTNAME_STRING(CUFAllocDescriptor), fctTy);1385 }1386 1387 mlir::Type structTy = typeConverter.convertBoxTypeAsStruct(boxTy);1388 std::size_t boxSize = dl->getTypeSizeInBits(structTy) / 8;1389 mlir::Value sizeInBytes =1390 fir::genConstantIndex(loc, llvmIntPtrType, rewriter, boxSize);1391 llvm::SmallVector args = {sizeInBytes, sourceFile, sourceLine};1392 return mlir::LLVM::CallOp::create(rewriter, loc, fctTy,1393 RTNAME_STRING(CUFAllocDescriptor), args)1394 .getResult();1395}1396 1397/// Get the address of the type descriptor global variable that was created by1398/// lowering for derived type \p recType.1399template <typename ModOpTy>1400static mlir::Value1401getTypeDescriptor(ModOpTy mod, mlir::ConversionPatternRewriter &rewriter,1402 mlir::Location loc, fir::RecordType recType,1403 const fir::FIRToLLVMPassOptions &options) {1404 std::string name =1405 options.typeDescriptorsRenamedForAssembly1406 ? fir::NameUniquer::getTypeDescriptorAssemblyName(recType.getName())1407 : fir::NameUniquer::getTypeDescriptorName(recType.getName());1408 mlir::Type llvmPtrTy = ::getLlvmPtrType(mod.getContext());1409 mlir::DataLayout dataLayout(mod);1410 if (auto global = mod.template lookupSymbol<fir::GlobalOp>(name))1411 return replaceWithAddrOfOrASCast(1412 rewriter, loc, fir::factory::getGlobalAddressSpace(&dataLayout),1413 fir::factory::getProgramAddressSpace(&dataLayout), global.getSymName(),1414 llvmPtrTy);1415 // The global may have already been translated to LLVM.1416 if (auto global = mod.template lookupSymbol<mlir::LLVM::GlobalOp>(name))1417 return replaceWithAddrOfOrASCast(1418 rewriter, loc, global.getAddrSpace(),1419 fir::factory::getProgramAddressSpace(&dataLayout), global.getSymName(),1420 llvmPtrTy);1421 // Type info derived types do not have type descriptors since they are the1422 // types defining type descriptors.1423 if (options.ignoreMissingTypeDescriptors ||1424 fir::NameUniquer::belongsToModule(1425 name, Fortran::semantics::typeInfoBuiltinModule))1426 return mlir::LLVM::ZeroOp::create(rewriter, loc, llvmPtrTy);1427 1428 if (!options.skipExternalRttiDefinition)1429 fir::emitFatalError(loc,1430 "runtime derived type info descriptor was not "1431 "generated and skipExternalRttiDefinition and "1432 "ignoreMissingTypeDescriptors options are not set");1433 1434 // Rtti for a derived type defined in another compilation unit and for which1435 // rtti was not defined in lowering because of the skipExternalRttiDefinition1436 // option. Generate the object declaration now.1437 auto insertPt = rewriter.saveInsertionPoint();1438 rewriter.setInsertionPoint(mod.getBody(), mod.getBody()->end());1439 mlir::LLVM::GlobalOp global = mlir::LLVM::GlobalOp::create(1440 rewriter, loc, llvmPtrTy, /*constant=*/true,1441 mlir::LLVM::Linkage::External, name, mlir::Attribute());1442 rewriter.restoreInsertionPoint(insertPt);1443 return mlir::LLVM::AddressOfOp::create(rewriter, loc, llvmPtrTy,1444 global.getSymName());1445}1446 1447/// Common base class for embox to descriptor conversion.1448template <typename OP>1449struct EmboxCommonConversion : public fir::FIROpConversion<OP> {1450 using fir::FIROpConversion<OP>::FIROpConversion;1451 using TypePair = typename fir::FIROpConversion<OP>::TypePair;1452 1453 static int getCFIAttr(fir::BaseBoxType boxTy) {1454 auto eleTy = boxTy.getEleTy();1455 if (mlir::isa<fir::PointerType>(eleTy))1456 return CFI_attribute_pointer;1457 if (mlir::isa<fir::HeapType>(eleTy))1458 return CFI_attribute_allocatable;1459 return CFI_attribute_other;1460 }1461 1462 mlir::Value getCharacterByteSize(mlir::Location loc,1463 mlir::ConversionPatternRewriter &rewriter,1464 fir::CharacterType charTy,1465 mlir::ValueRange lenParams) const {1466 auto i64Ty = mlir::IntegerType::get(rewriter.getContext(), 64);1467 mlir::Value size = genTypeStrideInBytes(1468 loc, i64Ty, rewriter, this->convertType(charTy), this->getDataLayout());1469 if (charTy.hasConstantLen())1470 return size; // Length accounted for in the genTypeStrideInBytes GEP.1471 // Otherwise, multiply the single character size by the length.1472 assert(!lenParams.empty());1473 auto len64 = fir::FIROpConversion<OP>::integerCast(loc, rewriter, i64Ty,1474 lenParams.back());1475 return mlir::LLVM::MulOp::create(rewriter, loc, i64Ty, size, len64);1476 }1477 1478 // Get the element size and CFI type code of the boxed value.1479 std::tuple<mlir::Value, mlir::Value> getSizeAndTypeCode(1480 mlir::Location loc, mlir::ConversionPatternRewriter &rewriter,1481 mlir::Type boxEleTy, mlir::ValueRange lenParams = {}) const {1482 const mlir::DataLayout &dataLayout = this->getDataLayout();1483 auto i64Ty = mlir::IntegerType::get(rewriter.getContext(), 64);1484 if (auto eleTy = fir::dyn_cast_ptrEleTy(boxEleTy))1485 boxEleTy = eleTy;1486 if (auto seqTy = mlir::dyn_cast<fir::SequenceType>(boxEleTy))1487 return getSizeAndTypeCode(loc, rewriter, seqTy.getEleTy(), lenParams);1488 if (mlir::isa<mlir::NoneType>(1489 boxEleTy)) // unlimited polymorphic or assumed type1490 return {mlir::LLVM::ConstantOp::create(rewriter, loc, i64Ty, 0),1491 this->genConstantOffset(loc, rewriter, CFI_type_other)};1492 mlir::Value typeCodeVal = this->genConstantOffset(1493 loc, rewriter,1494 fir::getTypeCode(boxEleTy, this->lowerTy().getKindMap()));1495 if (fir::isa_integer(boxEleTy) ||1496 mlir::dyn_cast<fir::LogicalType>(boxEleTy) || fir::isa_real(boxEleTy) ||1497 fir::isa_complex(boxEleTy))1498 return {genTypeStrideInBytes(loc, i64Ty, rewriter,1499 this->convertType(boxEleTy), dataLayout),1500 typeCodeVal};1501 if (auto charTy = mlir::dyn_cast<fir::CharacterType>(boxEleTy))1502 return {getCharacterByteSize(loc, rewriter, charTy, lenParams),1503 typeCodeVal};1504 if (fir::isa_ref_type(boxEleTy)) {1505 auto ptrTy = ::getLlvmPtrType(rewriter.getContext());1506 return {genTypeStrideInBytes(loc, i64Ty, rewriter, ptrTy, dataLayout),1507 typeCodeVal};1508 }1509 if (mlir::isa<fir::RecordType>(boxEleTy))1510 return {genTypeStrideInBytes(loc, i64Ty, rewriter,1511 this->convertType(boxEleTy), dataLayout),1512 typeCodeVal};1513 fir::emitFatalError(loc, "unhandled type in fir.box code generation");1514 }1515 1516 /// Basic pattern to write a field in the descriptor1517 mlir::Value insertField(mlir::ConversionPatternRewriter &rewriter,1518 mlir::Location loc, mlir::Value dest,1519 llvm::ArrayRef<std::int64_t> fldIndexes,1520 mlir::Value value, bool bitcast = false) const {1521 auto boxTy = dest.getType();1522 auto fldTy = this->getBoxEleTy(boxTy, fldIndexes);1523 if (!bitcast)1524 value = this->integerCast(loc, rewriter, fldTy, value);1525 // bitcast are no-ops with LLVM opaque pointers.1526 return mlir::LLVM::InsertValueOp::create(rewriter, loc, dest, value,1527 fldIndexes);1528 }1529 1530 inline mlir::Value1531 insertBaseAddress(mlir::ConversionPatternRewriter &rewriter,1532 mlir::Location loc, mlir::Value dest,1533 mlir::Value base) const {1534 return insertField(rewriter, loc, dest, {kAddrPosInBox}, base,1535 /*bitCast=*/true);1536 }1537 1538 inline mlir::Value insertLowerBound(mlir::ConversionPatternRewriter &rewriter,1539 mlir::Location loc, mlir::Value dest,1540 unsigned dim, mlir::Value lb) const {1541 return insertField(rewriter, loc, dest,1542 {kDimsPosInBox, dim, kDimLowerBoundPos}, lb);1543 }1544 1545 inline mlir::Value insertExtent(mlir::ConversionPatternRewriter &rewriter,1546 mlir::Location loc, mlir::Value dest,1547 unsigned dim, mlir::Value extent) const {1548 return insertField(rewriter, loc, dest, {kDimsPosInBox, dim, kDimExtentPos},1549 extent);1550 }1551 1552 inline mlir::Value insertStride(mlir::ConversionPatternRewriter &rewriter,1553 mlir::Location loc, mlir::Value dest,1554 unsigned dim, mlir::Value stride) const {1555 return insertField(rewriter, loc, dest, {kDimsPosInBox, dim, kDimStridePos},1556 stride);1557 }1558 1559 template <typename ModOpTy>1560 mlir::Value populateDescriptor(mlir::Location loc, ModOpTy mod,1561 fir::BaseBoxType boxTy, mlir::Type inputType,1562 mlir::ConversionPatternRewriter &rewriter,1563 unsigned rank, mlir::Value eleSize,1564 mlir::Value cfiTy, mlir::Value typeDesc,1565 int allocatorIdx = kDefaultAllocator,1566 mlir::Value extraField = {}) const {1567 auto llvmBoxTy = this->lowerTy().convertBoxTypeAsStruct(boxTy, rank);1568 bool isUnlimitedPolymorphic = fir::isUnlimitedPolymorphicType(boxTy);1569 bool useInputType = fir::isPolymorphicType(boxTy) || isUnlimitedPolymorphic;1570 mlir::Value descriptor =1571 mlir::LLVM::UndefOp::create(rewriter, loc, llvmBoxTy);1572 descriptor =1573 insertField(rewriter, loc, descriptor, {kElemLenPosInBox}, eleSize);1574 descriptor = insertField(rewriter, loc, descriptor, {kVersionPosInBox},1575 this->genI32Constant(loc, rewriter, CFI_VERSION));1576 descriptor = insertField(rewriter, loc, descriptor, {kRankPosInBox},1577 this->genI32Constant(loc, rewriter, rank));1578 descriptor = insertField(rewriter, loc, descriptor, {kTypePosInBox}, cfiTy);1579 descriptor =1580 insertField(rewriter, loc, descriptor, {kAttributePosInBox},1581 this->genI32Constant(loc, rewriter, getCFIAttr(boxTy)));1582 1583 const bool hasAddendum = fir::boxHasAddendum(boxTy);1584 1585 if (extraField) {1586 // Make sure to set the addendum presence flag according to the1587 // destination box.1588 if (hasAddendum) {1589 auto maskAttr = mlir::IntegerAttr::get(1590 rewriter.getIntegerType(8, /*isSigned=*/false),1591 llvm::APInt(8, (uint64_t)_CFI_ADDENDUM_FLAG, /*isSigned=*/false));1592 mlir::LLVM::ConstantOp mask = mlir::LLVM::ConstantOp::create(1593 rewriter, loc, rewriter.getI8Type(), maskAttr);1594 extraField = mlir::LLVM::OrOp::create(rewriter, loc, extraField, mask);1595 } else {1596 auto maskAttr = mlir::IntegerAttr::get(1597 rewriter.getIntegerType(8, /*isSigned=*/false),1598 llvm::APInt(8, (uint64_t)~_CFI_ADDENDUM_FLAG, /*isSigned=*/true));1599 mlir::LLVM::ConstantOp mask = mlir::LLVM::ConstantOp::create(1600 rewriter, loc, rewriter.getI8Type(), maskAttr);1601 extraField = mlir::LLVM::AndOp::create(rewriter, loc, extraField, mask);1602 }1603 // Extra field value is provided so just use it.1604 descriptor =1605 insertField(rewriter, loc, descriptor, {kExtraPosInBox}, extraField);1606 } else {1607 // Compute the value of the extra field based on allocator_idx and1608 // addendum present.1609 unsigned extra = allocatorIdx << _CFI_ALLOCATOR_IDX_SHIFT;1610 if (hasAddendum)1611 extra |= _CFI_ADDENDUM_FLAG;1612 descriptor = insertField(rewriter, loc, descriptor, {kExtraPosInBox},1613 this->genI32Constant(loc, rewriter, extra));1614 }1615 1616 if (hasAddendum) {1617 unsigned typeDescFieldId = getTypeDescFieldId(boxTy);1618 if (!typeDesc) {1619 if (useInputType) {1620 mlir::Type innerType = fir::unwrapInnerType(inputType);1621 if (innerType && mlir::isa<fir::RecordType>(innerType)) {1622 auto recTy = mlir::dyn_cast<fir::RecordType>(innerType);1623 typeDesc =1624 getTypeDescriptor(mod, rewriter, loc, recTy, this->options);1625 } else {1626 // Unlimited polymorphic type descriptor with no record type. Set1627 // type descriptor address to a clean state.1628 typeDesc = mlir::LLVM::ZeroOp::create(1629 rewriter, loc, ::getLlvmPtrType(mod.getContext()));1630 }1631 } else {1632 typeDesc = getTypeDescriptor(1633 mod, rewriter, loc, fir::unwrapIfDerived(boxTy), this->options);1634 }1635 }1636 if (typeDesc)1637 descriptor =1638 insertField(rewriter, loc, descriptor, {typeDescFieldId}, typeDesc,1639 /*bitCast=*/true);1640 // Always initialize the length parameter field to zero to avoid issues1641 // with uninitialized values in Fortran code trying to compare physical1642 // representation of derived types with pointer/allocatable components.1643 // This has been seen in hashing algorithms using TRANSFER.1644 mlir::Value zero =1645 fir::genConstantIndex(loc, rewriter.getI64Type(), rewriter, 0);1646 descriptor = insertField(rewriter, loc, descriptor,1647 {getLenParamFieldId(boxTy), 0}, zero);1648 }1649 return descriptor;1650 }1651 1652 // Template used for fir::EmboxOp and fir::cg::XEmboxOp1653 template <typename BOX>1654 std::tuple<fir::BaseBoxType, mlir::Value, mlir::Value>1655 consDescriptorPrefix(BOX box, mlir::Type inputType,1656 mlir::ConversionPatternRewriter &rewriter, unsigned rank,1657 [[maybe_unused]] mlir::ValueRange substrParams,1658 mlir::ValueRange lenParams, mlir::Value sourceBox = {},1659 mlir::Type sourceBoxType = {}) const {1660 auto loc = box.getLoc();1661 auto boxTy = mlir::dyn_cast<fir::BaseBoxType>(box.getType());1662 bool useInputType = fir::isPolymorphicType(boxTy) &&1663 !fir::isUnlimitedPolymorphicType(inputType);1664 llvm::SmallVector<mlir::Value> typeparams = lenParams;1665 if constexpr (!std::is_same_v<BOX, fir::EmboxOp>) {1666 if (!box.getSubstr().empty() && fir::hasDynamicSize(boxTy.getEleTy()))1667 typeparams.push_back(substrParams[1]);1668 }1669 1670 int allocatorIdx = 0;1671 if constexpr (std::is_same_v<BOX, fir::EmboxOp> ||1672 std::is_same_v<BOX, fir::cg::XEmboxOp>) {1673 if (box.getAllocatorIdx())1674 allocatorIdx = *box.getAllocatorIdx();1675 }1676 1677 // Write each of the fields with the appropriate values.1678 // When emboxing an element to a polymorphic descriptor, use the1679 // input type since the destination descriptor type has not the exact1680 // information.1681 auto [eleSize, cfiTy] = getSizeAndTypeCode(1682 loc, rewriter, useInputType ? inputType : boxTy.getEleTy(), typeparams);1683 1684 mlir::Value typeDesc;1685 mlir::Value extraField;1686 // When emboxing to a polymorphic box, get the type descriptor, type code1687 // and element size from the source box if any.1688 if (fir::isPolymorphicType(boxTy) && sourceBox) {1689 TypePair sourceBoxTyPair = this->getBoxTypePair(sourceBoxType);1690 typeDesc =1691 this->loadTypeDescAddress(loc, sourceBoxTyPair, sourceBox, rewriter);1692 mlir::Type idxTy = this->lowerTy().indexType();1693 eleSize = this->getElementSizeFromBox(loc, idxTy, sourceBoxTyPair,1694 sourceBox, rewriter);1695 cfiTy = this->getValueFromBox(loc, sourceBoxTyPair, sourceBox,1696 cfiTy.getType(), rewriter, kTypePosInBox);1697 extraField =1698 this->getExtraFromBox(loc, sourceBoxTyPair, sourceBox, rewriter);1699 }1700 1701 mlir::Value descriptor;1702 if (auto gpuMod = box->template getParentOfType<mlir::gpu::GPUModuleOp>())1703 descriptor = populateDescriptor(loc, gpuMod, boxTy, inputType, rewriter,1704 rank, eleSize, cfiTy, typeDesc,1705 allocatorIdx, extraField);1706 else if (auto mod = box->template getParentOfType<mlir::ModuleOp>())1707 descriptor = populateDescriptor(loc, mod, boxTy, inputType, rewriter,1708 rank, eleSize, cfiTy, typeDesc,1709 allocatorIdx, extraField);1710 1711 return {boxTy, descriptor, eleSize};1712 }1713 1714 std::tuple<fir::BaseBoxType, mlir::Value, mlir::Value>1715 consDescriptorPrefix(fir::cg::XReboxOp box, mlir::Value loweredBox,1716 mlir::ConversionPatternRewriter &rewriter, unsigned rank,1717 mlir::ValueRange substrParams,1718 mlir::ValueRange lenParams,1719 mlir::Value typeDesc = {}) const {1720 auto loc = box.getLoc();1721 auto boxTy = mlir::dyn_cast<fir::BaseBoxType>(box.getType());1722 auto inputBoxTy = mlir::dyn_cast<fir::BaseBoxType>(box.getBox().getType());1723 auto inputBoxTyPair = this->getBoxTypePair(inputBoxTy);1724 llvm::SmallVector<mlir::Value> typeparams = lenParams;1725 if (!box.getSubstr().empty() && fir::hasDynamicSize(boxTy.getEleTy()))1726 typeparams.push_back(substrParams[1]);1727 1728 auto [eleSize, cfiTy] =1729 getSizeAndTypeCode(loc, rewriter, boxTy.getEleTy(), typeparams);1730 1731 // Reboxing to a polymorphic entity. eleSize and type code need to1732 // be retrieved from the initial box and propagated to the new box.1733 // If the initial box has an addendum, the type desc must be propagated as1734 // well.1735 if (fir::isPolymorphicType(boxTy)) {1736 mlir::Type idxTy = this->lowerTy().indexType();1737 eleSize = this->getElementSizeFromBox(loc, idxTy, inputBoxTyPair,1738 loweredBox, rewriter);1739 cfiTy = this->getValueFromBox(loc, inputBoxTyPair, loweredBox,1740 cfiTy.getType(), rewriter, kTypePosInBox);1741 // TODO: For initial box that are unlimited polymorphic entities, this1742 // code must be made conditional because unlimited polymorphic entities1743 // with intrinsic type spec does not have addendum.1744 if (fir::boxHasAddendum(inputBoxTy))1745 typeDesc = this->loadTypeDescAddress(loc, inputBoxTyPair, loweredBox,1746 rewriter);1747 }1748 1749 mlir::Value extraField =1750 this->getExtraFromBox(loc, inputBoxTyPair, loweredBox, rewriter);1751 1752 mlir::Value descriptor;1753 if (auto gpuMod = box->template getParentOfType<mlir::gpu::GPUModuleOp>())1754 descriptor =1755 populateDescriptor(loc, gpuMod, boxTy, box.getBox().getType(),1756 rewriter, rank, eleSize, cfiTy, typeDesc,1757 /*allocatorIdx=*/kDefaultAllocator, extraField);1758 else if (auto mod = box->template getParentOfType<mlir::ModuleOp>())1759 descriptor =1760 populateDescriptor(loc, mod, boxTy, box.getBox().getType(), rewriter,1761 rank, eleSize, cfiTy, typeDesc,1762 /*allocatorIdx=*/kDefaultAllocator, extraField);1763 1764 return {boxTy, descriptor, eleSize};1765 }1766 1767 // Compute the base address of a fir.box given the indices from the slice.1768 // The indices from the "outer" dimensions (every dimension after the first1769 // one (included) that is not a compile time constant) must have been1770 // multiplied with the related extents and added together into \p outerOffset.1771 mlir::Value1772 genBoxOffsetGep(mlir::ConversionPatternRewriter &rewriter, mlir::Location loc,1773 mlir::Value base, mlir::Type llvmBaseObjectType,1774 mlir::Value outerOffset, mlir::ValueRange cstInteriorIndices,1775 mlir::ValueRange componentIndices,1776 std::optional<mlir::Value> substringOffset) const {1777 llvm::SmallVector<mlir::LLVM::GEPArg> gepArgs{outerOffset};1778 mlir::Type resultTy = llvmBaseObjectType;1779 // Fortran is column major, llvm GEP is row major: reverse the indices here.1780 for (mlir::Value interiorIndex : llvm::reverse(cstInteriorIndices)) {1781 auto arrayTy = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(resultTy);1782 if (!arrayTy)1783 fir::emitFatalError(1784 loc,1785 "corrupted GEP generated being generated in fir.embox/fir.rebox");1786 resultTy = arrayTy.getElementType();1787 gepArgs.push_back(interiorIndex);1788 }1789 llvm::SmallVector<mlir::Value> gepIndices =1790 convertSubcomponentIndices(loc, resultTy, componentIndices, &resultTy);1791 gepArgs.append(gepIndices.begin(), gepIndices.end());1792 if (substringOffset) {1793 if (auto arrayTy = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(resultTy)) {1794 gepArgs.push_back(*substringOffset);1795 resultTy = arrayTy.getElementType();1796 } else {1797 // If the CHARACTER length is dynamic, the whole base type should have1798 // degenerated to an llvm.ptr<i[width]>, and there should not be any1799 // cstInteriorIndices/componentIndices. The substring offset can be1800 // added to the outterOffset since it applies on the same LLVM type.1801 if (gepArgs.size() != 1)1802 fir::emitFatalError(loc,1803 "corrupted substring GEP in fir.embox/fir.rebox");1804 mlir::Type outterOffsetTy =1805 llvm::cast<mlir::Value>(gepArgs[0]).getType();1806 mlir::Value cast =1807 this->integerCast(loc, rewriter, outterOffsetTy, *substringOffset);1808 1809 gepArgs[0] = mlir::LLVM::AddOp::create(1810 rewriter, loc, outterOffsetTy, llvm::cast<mlir::Value>(gepArgs[0]),1811 cast);1812 }1813 }1814 mlir::Type llvmPtrTy = ::getLlvmPtrType(resultTy.getContext());1815 return mlir::LLVM::GEPOp::create(rewriter, loc, llvmPtrTy,1816 llvmBaseObjectType, base, gepArgs);1817 }1818 1819 template <typename BOX>1820 void1821 getSubcomponentIndices(BOX xbox, mlir::Value memref,1822 mlir::ValueRange operands,1823 mlir::SmallVectorImpl<mlir::Value> &indices) const {1824 // For each field in the path add the offset to base via the args list.1825 // In the most general case, some offsets must be computed since1826 // they are not be known until runtime.1827 if (fir::hasDynamicSize(fir::unwrapSequenceType(1828 fir::unwrapPassByRefType(memref.getType()))))1829 TODO(xbox.getLoc(),1830 "fir.embox codegen dynamic size component in derived type");1831 indices.append(operands.begin() + xbox.getSubcomponentOperandIndex(),1832 operands.begin() + xbox.getSubcomponentOperandIndex() +1833 xbox.getSubcomponent().size());1834 }1835 1836 /// If the embox is not in a globalOp body, allocate storage for the box;1837 /// store the value inside and return the generated alloca. Return the input1838 /// value otherwise.1839 mlir::Value1840 placeInMemoryIfNotGlobalInit(mlir::ConversionPatternRewriter &rewriter,1841 mlir::Location loc, mlir::Type boxTy,1842 mlir::Value boxValue,1843 bool needDeviceAllocation = false) const {1844 if (isInGlobalOp(rewriter))1845 return boxValue;1846 mlir::Type llvmBoxTy = boxValue.getType();1847 mlir::Value storage;1848 if (needDeviceAllocation) {1849 auto mod = boxValue.getDefiningOp()->getParentOfType<mlir::ModuleOp>();1850 auto baseBoxTy = mlir::dyn_cast<fir::BaseBoxType>(boxTy);1851 storage =1852 genCUFAllocDescriptor(loc, rewriter, mod, baseBoxTy, this->lowerTy());1853 } else {1854 storage = this->genAllocaAndAddrCastWithType(loc, llvmBoxTy, defaultAlign,1855 rewriter);1856 }1857 auto storeOp =1858 mlir::LLVM::StoreOp::create(rewriter, loc, boxValue, storage);1859 this->attachTBAATag(storeOp, boxTy, boxTy, nullptr);1860 return storage;1861 }1862 1863 /// Compute the extent of a triplet slice (lb:ub:step).1864 mlir::Value computeTripletExtent(mlir::ConversionPatternRewriter &rewriter,1865 mlir::Location loc, mlir::Value lb,1866 mlir::Value ub, mlir::Value step,1867 mlir::Value zero, mlir::Type type) const {1868 lb = this->integerCast(loc, rewriter, type, lb);1869 ub = this->integerCast(loc, rewriter, type, ub);1870 step = this->integerCast(loc, rewriter, type, step);1871 zero = this->integerCast(loc, rewriter, type, zero);1872 mlir::Value extent = mlir::LLVM::SubOp::create(rewriter, loc, type, ub, lb);1873 extent = mlir::LLVM::AddOp::create(rewriter, loc, type, extent, step);1874 extent = mlir::LLVM::SDivOp::create(rewriter, loc, type, extent, step);1875 // If the resulting extent is negative (`ub-lb` and `step` have different1876 // signs), zero must be returned instead.1877 auto cmp = mlir::LLVM::ICmpOp::create(1878 rewriter, loc, mlir::LLVM::ICmpPredicate::sgt, extent, zero);1879 return mlir::LLVM::SelectOp::create(rewriter, loc, cmp, extent, zero);1880 }1881};1882 1883/// Create a generic box on a memory reference. This conversions lowers the1884/// abstract box to the appropriate, initialized descriptor.1885struct EmboxOpConversion : public EmboxCommonConversion<fir::EmboxOp> {1886 using EmboxCommonConversion::EmboxCommonConversion;1887 1888 llvm::LogicalResult1889 matchAndRewrite(fir::EmboxOp embox, OpAdaptor adaptor,1890 mlir::ConversionPatternRewriter &rewriter) const override {1891 mlir::ValueRange operands = adaptor.getOperands();1892 mlir::Value sourceBox;1893 mlir::Type sourceBoxType;1894 if (embox.getSourceBox()) {1895 sourceBox = operands[embox.getSourceBoxOperandIndex()];1896 sourceBoxType = embox.getSourceBox().getType();1897 }1898 assert(!embox.getShape() && "There should be no dims on this embox op");1899 auto [boxTy, dest, eleSize] = consDescriptorPrefix(1900 embox, fir::unwrapRefType(embox.getMemref().getType()), rewriter,1901 /*rank=*/0, /*substrParams=*/mlir::ValueRange{},1902 adaptor.getTypeparams(), sourceBox, sourceBoxType);1903 dest = insertBaseAddress(rewriter, embox.getLoc(), dest, operands[0]);1904 if (fir::isDerivedTypeWithLenParams(boxTy)) {1905 TODO(embox.getLoc(),1906 "fir.embox codegen of derived with length parameters");1907 return mlir::failure();1908 }1909 auto result =1910 placeInMemoryIfNotGlobalInit(rewriter, embox.getLoc(), boxTy, dest);1911 rewriter.replaceOp(embox, result);1912 return mlir::success();1913 }1914};1915 1916static bool isDeviceAllocation(mlir::Value val, mlir::Value adaptorVal) {1917 if (val.getDefiningOp() &&1918 val.getDefiningOp()->getParentOfType<mlir::gpu::GPUModuleOp>())1919 return false;1920 // Check if the global symbol is in the device module.1921 if (auto addr = mlir::dyn_cast_or_null<fir::AddrOfOp>(val.getDefiningOp()))1922 if (auto gpuMod =1923 addr->getParentOfType<mlir::ModuleOp>()1924 .lookupSymbol<mlir::gpu::GPUModuleOp>(cudaDeviceModuleName))1925 if (gpuMod.lookupSymbol<mlir::LLVM::GlobalOp>(addr.getSymbol()) ||1926 gpuMod.lookupSymbol<fir::GlobalOp>(addr.getSymbol()))1927 return true;1928 1929 if (auto loadOp = mlir::dyn_cast_or_null<fir::LoadOp>(val.getDefiningOp()))1930 return isDeviceAllocation(loadOp.getMemref(), {});1931 if (auto boxAddrOp =1932 mlir::dyn_cast_or_null<fir::BoxAddrOp>(val.getDefiningOp()))1933 return isDeviceAllocation(boxAddrOp.getVal(), {});1934 if (auto convertOp =1935 mlir::dyn_cast_or_null<fir::ConvertOp>(val.getDefiningOp()))1936 return isDeviceAllocation(convertOp.getValue(), {});1937 if (!val.getDefiningOp() && adaptorVal) {1938 if (auto blockArg = llvm::cast<mlir::BlockArgument>(adaptorVal)) {1939 if (blockArg.getOwner() && blockArg.getOwner()->getParentOp() &&1940 blockArg.getOwner()->isEntryBlock()) {1941 if (auto func = mlir::dyn_cast_or_null<mlir::FunctionOpInterface>(1942 *blockArg.getOwner()->getParentOp())) {1943 auto argAttrs = func.getArgAttrs(blockArg.getArgNumber());1944 for (auto attr : argAttrs) {1945 if (attr.getName().getValue().ends_with(cuf::getDataAttrName())) {1946 auto dataAttr =1947 mlir::dyn_cast<cuf::DataAttributeAttr>(attr.getValue());1948 if (dataAttr.getValue() != cuf::DataAttribute::Pinned &&1949 dataAttr.getValue() != cuf::DataAttribute::Unified)1950 return true;1951 }1952 }1953 }1954 }1955 }1956 }1957 if (auto callOp = mlir::dyn_cast_or_null<fir::CallOp>(val.getDefiningOp()))1958 if (callOp.getCallee() &&1959 (callOp.getCallee().value().getRootReference().getValue().starts_with(1960 RTNAME_STRING(CUFMemAlloc)) ||1961 callOp.getCallee().value().getRootReference().getValue().starts_with(1962 RTNAME_STRING(CUFAllocDescriptor)) ||1963 callOp.getCallee().value().getRootReference().getValue() ==1964 "__tgt_acc_get_deviceptr"))1965 return true;1966 return false;1967}1968 1969/// Create a generic box on a memory reference.1970struct XEmboxOpConversion : public EmboxCommonConversion<fir::cg::XEmboxOp> {1971 using EmboxCommonConversion::EmboxCommonConversion;1972 1973 llvm::LogicalResult1974 matchAndRewrite(fir::cg::XEmboxOp xbox, OpAdaptor adaptor,1975 mlir::ConversionPatternRewriter &rewriter) const override {1976 mlir::ValueRange operands = adaptor.getOperands();1977 mlir::Value sourceBox;1978 mlir::Type sourceBoxType;1979 if (xbox.getSourceBox()) {1980 sourceBox = operands[xbox.getSourceBoxOperandIndex()];1981 sourceBoxType = xbox.getSourceBox().getType();1982 }1983 auto [boxTy, dest, resultEleSize] = consDescriptorPrefix(1984 xbox, fir::unwrapRefType(xbox.getMemref().getType()), rewriter,1985 xbox.getOutRank(), adaptor.getSubstr(), adaptor.getLenParams(),1986 sourceBox, sourceBoxType);1987 // Generate the triples in the dims field of the descriptor1988 auto i64Ty = mlir::IntegerType::get(xbox.getContext(), 64);1989 assert(!xbox.getShape().empty() && "must have a shape");1990 unsigned shapeOffset = xbox.getShapeOperandIndex();1991 bool hasShift = !xbox.getShift().empty();1992 unsigned shiftOffset = xbox.getShiftOperandIndex();1993 bool hasSlice = !xbox.getSlice().empty();1994 unsigned sliceOffset = xbox.getSliceOperandIndex();1995 mlir::Location loc = xbox.getLoc();1996 mlir::Value zero = fir::genConstantIndex(loc, i64Ty, rewriter, 0);1997 mlir::Value one = fir::genConstantIndex(loc, i64Ty, rewriter, 1);1998 mlir::Value prevPtrOff = one;1999 mlir::Type eleTy = boxTy.getEleTy();2000 const unsigned rank = xbox.getRank();2001 llvm::SmallVector<mlir::Value> cstInteriorIndices;2002 unsigned constRows = 0;2003 mlir::Value ptrOffset = zero;2004 mlir::Type memEleTy = fir::dyn_cast_ptrEleTy(xbox.getMemref().getType());2005 assert(mlir::isa<fir::SequenceType>(memEleTy));2006 auto seqTy = mlir::cast<fir::SequenceType>(memEleTy);2007 mlir::Type seqEleTy = seqTy.getEleTy();2008 // Adjust the element scaling factor if the element is a dependent type.2009 if (fir::hasDynamicSize(seqEleTy)) {2010 if (auto charTy = mlir::dyn_cast<fir::CharacterType>(seqEleTy)) {2011 // The GEP pointer type decays to llvm.ptr<i[width]>.2012 // The scaling factor is the runtime value of the length.2013 assert(!adaptor.getLenParams().empty());2014 prevPtrOff = FIROpConversion::integerCast(2015 loc, rewriter, i64Ty, adaptor.getLenParams().back());2016 } else if (mlir::isa<fir::RecordType>(seqEleTy)) {2017 // prevPtrOff = ;2018 TODO(loc, "generate call to calculate size of PDT");2019 } else {2020 fir::emitFatalError(loc, "unexpected dynamic type");2021 }2022 } else {2023 constRows = seqTy.getConstantRows();2024 }2025 2026 const auto hasSubcomp = !xbox.getSubcomponent().empty();2027 const bool hasSubstr = !xbox.getSubstr().empty();2028 // Initial element stride that will be use to compute the step in2029 // each dimension. Initially, this is the size of the input element.2030 // Note that when there are no components/substring, the resultEleSize2031 // that was previously computed matches the input element size.2032 mlir::Value prevDimByteStride = resultEleSize;2033 if (hasSubcomp) {2034 // We have a subcomponent. The step value needs to be the number of2035 // bytes per element (which is a derived type).2036 prevDimByteStride = genTypeStrideInBytes(2037 loc, i64Ty, rewriter, convertType(seqEleTy), getDataLayout());2038 } else if (hasSubstr) {2039 // We have a substring. The step value needs to be the number of bytes2040 // per CHARACTER element.2041 auto charTy = mlir::cast<fir::CharacterType>(seqEleTy);2042 if (fir::hasDynamicSize(charTy)) {2043 prevDimByteStride =2044 getCharacterByteSize(loc, rewriter, charTy, adaptor.getLenParams());2045 } else {2046 prevDimByteStride = fir::genConstantIndex(2047 loc, i64Ty, rewriter,2048 charTy.getLen() * lowerTy().characterBitsize(charTy) / 8);2049 }2050 }2051 2052 // Process the array subspace arguments (shape, shift, etc.), if any,2053 // translating everything to values in the descriptor wherever the entity2054 // has a dynamic array dimension.2055 for (unsigned di = 0, descIdx = 0; di < rank; ++di) {2056 mlir::Value extent =2057 integerCast(loc, rewriter, i64Ty, operands[shapeOffset]);2058 mlir::Value outerExtent = extent;2059 bool skipNext = false;2060 if (hasSlice) {2061 mlir::Value off =2062 integerCast(loc, rewriter, i64Ty, operands[sliceOffset]);2063 mlir::Value adj = one;2064 if (hasShift)2065 adj = integerCast(loc, rewriter, i64Ty, operands[shiftOffset]);2066 auto ao = mlir::LLVM::SubOp::create(rewriter, loc, i64Ty, off, adj);2067 if (constRows > 0) {2068 cstInteriorIndices.push_back(ao);2069 } else {2070 auto dimOff =2071 mlir::LLVM::MulOp::create(rewriter, loc, i64Ty, ao, prevPtrOff);2072 ptrOffset = mlir::LLVM::AddOp::create(rewriter, loc, i64Ty, dimOff,2073 ptrOffset);2074 }2075 if (mlir::isa_and_nonnull<fir::UndefOp>(2076 xbox.getSlice()[3 * di + 1].getDefiningOp())) {2077 // This dimension contains a scalar expression in the array slice op.2078 // The dimension is loop invariant, will be dropped, and will not2079 // appear in the descriptor.2080 skipNext = true;2081 }2082 }2083 if (!skipNext) {2084 // store extent2085 if (hasSlice)2086 extent = computeTripletExtent(rewriter, loc, operands[sliceOffset],2087 operands[sliceOffset + 1],2088 operands[sliceOffset + 2], zero, i64Ty);2089 // Lower bound is normalized to 0 for BIND(C) interoperability.2090 mlir::Value lb = zero;2091 const bool isaPointerOrAllocatable =2092 mlir::isa<fir::PointerType, fir::HeapType>(eleTy);2093 // Lower bound is defaults to 1 for POINTER, ALLOCATABLE, and2094 // denormalized descriptors.2095 if (isaPointerOrAllocatable || !normalizedLowerBound(xbox))2096 lb = one;2097 // If there is a shifted origin, and no fir.slice, and this is not2098 // a normalized descriptor then use the value from the shift op as2099 // the lower bound.2100 if (hasShift && !(hasSlice || hasSubcomp || hasSubstr) &&2101 (isaPointerOrAllocatable || !normalizedLowerBound(xbox))) {2102 lb = integerCast(loc, rewriter, i64Ty, operands[shiftOffset]);2103 auto extentIsEmpty = mlir::LLVM::ICmpOp::create(2104 rewriter, loc, mlir::LLVM::ICmpPredicate::eq, extent, zero);2105 lb = mlir::LLVM::SelectOp::create(rewriter, loc, extentIsEmpty, one,2106 lb);2107 }2108 dest = insertLowerBound(rewriter, loc, dest, descIdx, lb);2109 2110 dest = insertExtent(rewriter, loc, dest, descIdx, extent);2111 2112 // store step (scaled by shaped extent)2113 mlir::Value step = prevDimByteStride;2114 if (hasSlice) {2115 mlir::Value sliceStep =2116 integerCast(loc, rewriter, i64Ty, operands[sliceOffset + 2]);2117 step =2118 mlir::LLVM::MulOp::create(rewriter, loc, i64Ty, step, sliceStep);2119 }2120 dest = insertStride(rewriter, loc, dest, descIdx, step);2121 ++descIdx;2122 }2123 2124 // compute the stride and offset for the next natural dimension2125 prevDimByteStride = mlir::LLVM::MulOp::create(2126 rewriter, loc, i64Ty, prevDimByteStride, outerExtent);2127 if (constRows == 0)2128 prevPtrOff = mlir::LLVM::MulOp::create(rewriter, loc, i64Ty, prevPtrOff,2129 outerExtent);2130 else2131 --constRows;2132 2133 // increment iterators2134 ++shapeOffset;2135 if (hasShift)2136 ++shiftOffset;2137 if (hasSlice)2138 sliceOffset += 3;2139 }2140 mlir::Value base = adaptor.getMemref();2141 if (hasSlice || hasSubcomp || hasSubstr) {2142 // Shift the base address.2143 llvm::SmallVector<mlir::Value> fieldIndices;2144 std::optional<mlir::Value> substringOffset;2145 if (hasSubcomp)2146 getSubcomponentIndices(xbox, xbox.getMemref(), operands, fieldIndices);2147 if (hasSubstr)2148 substringOffset = operands[xbox.getSubstrOperandIndex()];2149 mlir::Type llvmBaseType =2150 convertType(fir::unwrapRefType(xbox.getMemref().getType()));2151 base = genBoxOffsetGep(rewriter, loc, base, llvmBaseType, ptrOffset,2152 cstInteriorIndices, fieldIndices, substringOffset);2153 }2154 dest = insertBaseAddress(rewriter, loc, dest, base);2155 if (fir::isDerivedTypeWithLenParams(boxTy))2156 TODO(loc, "fir.embox codegen of derived with length parameters");2157 mlir::Value result = placeInMemoryIfNotGlobalInit(2158 rewriter, loc, boxTy, dest,2159 isDeviceAllocation(xbox.getMemref(), adaptor.getMemref()));2160 rewriter.replaceOp(xbox, result);2161 return mlir::success();2162 }2163 2164 /// Return true if `xbox` has a normalized lower bounds attribute. A box value2165 /// that is neither a POINTER nor an ALLOCATABLE should be normalized to a2166 /// zero origin lower bound for interoperability with BIND(C).2167 inline static bool normalizedLowerBound(fir::cg::XEmboxOp xbox) {2168 return xbox->hasAttr(fir::getNormalizedLowerBoundAttrName());2169 }2170};2171 2172/// Create a new box given a box reference.2173struct XReboxOpConversion : public EmboxCommonConversion<fir::cg::XReboxOp> {2174 using EmboxCommonConversion::EmboxCommonConversion;2175 2176 llvm::LogicalResult2177 matchAndRewrite(fir::cg::XReboxOp rebox, OpAdaptor adaptor,2178 mlir::ConversionPatternRewriter &rewriter) const override {2179 mlir::Location loc = rebox.getLoc();2180 mlir::Type idxTy = lowerTy().indexType();2181 mlir::Value loweredBox =2182 fixBoxInputInsideGlobalOp(rewriter, adaptor.getBox());2183 mlir::ValueRange operands = adaptor.getOperands();2184 2185 TypePair inputBoxTyPair = getBoxTypePair(rebox.getBox().getType());2186 2187 // Create new descriptor and fill its non-shape related data.2188 llvm::SmallVector<mlir::Value, 2> lenParams;2189 mlir::Type inputEleTy = getInputEleTy(rebox);2190 if (auto charTy = mlir::dyn_cast<fir::CharacterType>(inputEleTy)) {2191 if (charTy.hasConstantLen()) {2192 mlir::Value len =2193 fir::genConstantIndex(loc, idxTy, rewriter, charTy.getLen());2194 lenParams.emplace_back(len);2195 } else {2196 mlir::Value len = getElementSizeFromBox(loc, idxTy, inputBoxTyPair,2197 loweredBox, rewriter);2198 if (charTy.getFKind() != 1) {2199 assert(!isInGlobalOp(rewriter) &&2200 "character target in global op must have constant length");2201 mlir::Value width =2202 fir::genConstantIndex(loc, idxTy, rewriter, charTy.getFKind());2203 len = mlir::LLVM::SDivOp::create(rewriter, loc, idxTy, len, width);2204 }2205 lenParams.emplace_back(len);2206 }2207 } else if (auto recTy = mlir::dyn_cast<fir::RecordType>(inputEleTy)) {2208 if (recTy.getNumLenParams() != 0)2209 TODO(loc, "reboxing descriptor of derived type with length parameters");2210 }2211 2212 // Rebox on polymorphic entities needs to carry over the dynamic type.2213 mlir::Value typeDescAddr;2214 if (mlir::isa<fir::ClassType>(inputBoxTyPair.fir) &&2215 mlir::isa<fir::ClassType>(rebox.getType()))2216 typeDescAddr =2217 loadTypeDescAddress(loc, inputBoxTyPair, loweredBox, rewriter);2218 2219 auto [boxTy, dest, eleSize] =2220 consDescriptorPrefix(rebox, loweredBox, rewriter, rebox.getOutRank(),2221 adaptor.getSubstr(), lenParams, typeDescAddr);2222 2223 // Read input extents, strides, and base address2224 llvm::SmallVector<mlir::Value> inputExtents;2225 llvm::SmallVector<mlir::Value> inputStrides;2226 const unsigned inputRank = rebox.getRank();2227 for (unsigned dim = 0; dim < inputRank; ++dim) {2228 llvm::SmallVector<mlir::Value, 3> dimInfo =2229 getDimsFromBox(loc, {idxTy, idxTy, idxTy}, inputBoxTyPair, loweredBox,2230 dim, rewriter);2231 inputExtents.emplace_back(dimInfo[1]);2232 inputStrides.emplace_back(dimInfo[2]);2233 }2234 2235 mlir::Value baseAddr =2236 getBaseAddrFromBox(loc, inputBoxTyPair, loweredBox, rewriter);2237 2238 if (!rebox.getSlice().empty() || !rebox.getSubcomponent().empty())2239 return sliceBox(rebox, adaptor, boxTy, dest, baseAddr, inputExtents,2240 inputStrides, operands, rewriter);2241 return reshapeBox(rebox, adaptor, boxTy, dest, baseAddr, inputExtents,2242 inputStrides, operands, rewriter);2243 }2244 2245private:2246 /// Write resulting shape and base address in descriptor, and replace rebox2247 /// op.2248 llvm::LogicalResult2249 finalizeRebox(fir::cg::XReboxOp rebox, OpAdaptor adaptor,2250 mlir::Type destBoxTy, mlir::Value dest, mlir::Value base,2251 mlir::ValueRange lbounds, mlir::ValueRange extents,2252 mlir::ValueRange strides,2253 mlir::ConversionPatternRewriter &rewriter) const {2254 mlir::Location loc = rebox.getLoc();2255 mlir::Value zero =2256 fir::genConstantIndex(loc, lowerTy().indexType(), rewriter, 0);2257 mlir::Value one =2258 fir::genConstantIndex(loc, lowerTy().indexType(), rewriter, 1);2259 for (auto iter : llvm::enumerate(llvm::zip(extents, strides))) {2260 mlir::Value extent = std::get<0>(iter.value());2261 unsigned dim = iter.index();2262 mlir::Value lb = one;2263 if (!lbounds.empty()) {2264 lb = integerCast(loc, rewriter, lowerTy().indexType(), lbounds[dim]);2265 auto extentIsEmpty = mlir::LLVM::ICmpOp::create(2266 rewriter, loc, mlir::LLVM::ICmpPredicate::eq, extent, zero);2267 lb =2268 mlir::LLVM::SelectOp::create(rewriter, loc, extentIsEmpty, one, lb);2269 };2270 dest = insertLowerBound(rewriter, loc, dest, dim, lb);2271 dest = insertExtent(rewriter, loc, dest, dim, extent);2272 dest = insertStride(rewriter, loc, dest, dim, std::get<1>(iter.value()));2273 }2274 dest = insertBaseAddress(rewriter, loc, dest, base);2275 mlir::Value result = placeInMemoryIfNotGlobalInit(2276 rewriter, rebox.getLoc(), destBoxTy, dest,2277 isDeviceAllocation(rebox.getBox(), adaptor.getBox()));2278 rewriter.replaceOp(rebox, result);2279 return mlir::success();2280 }2281 2282 // Apply slice given the base address, extents and strides of the input box.2283 llvm::LogicalResult2284 sliceBox(fir::cg::XReboxOp rebox, OpAdaptor adaptor, mlir::Type destBoxTy,2285 mlir::Value dest, mlir::Value base, mlir::ValueRange inputExtents,2286 mlir::ValueRange inputStrides, mlir::ValueRange operands,2287 mlir::ConversionPatternRewriter &rewriter) const {2288 mlir::Location loc = rebox.getLoc();2289 mlir::Type byteTy = ::getI8Type(rebox.getContext());2290 mlir::Type idxTy = lowerTy().indexType();2291 mlir::Value zero = fir::genConstantIndex(loc, idxTy, rewriter, 0);2292 // Apply subcomponent and substring shift on base address.2293 if (!rebox.getSubcomponent().empty() || !rebox.getSubstr().empty()) {2294 // Cast to inputEleTy* so that a GEP can be used.2295 mlir::Type inputEleTy = getInputEleTy(rebox);2296 mlir::Type llvmBaseObjectType = convertType(inputEleTy);2297 llvm::SmallVector<mlir::Value> fieldIndices;2298 std::optional<mlir::Value> substringOffset;2299 if (!rebox.getSubcomponent().empty())2300 getSubcomponentIndices(rebox, rebox.getBox(), operands, fieldIndices);2301 if (!rebox.getSubstr().empty())2302 substringOffset = operands[rebox.getSubstrOperandIndex()];2303 base = genBoxOffsetGep(rewriter, loc, base, llvmBaseObjectType, zero,2304 /*cstInteriorIndices=*/{}, fieldIndices,2305 substringOffset);2306 }2307 2308 if (rebox.getSlice().empty())2309 // The array section is of the form array[%component][substring], keep2310 // the input array extents and strides.2311 return finalizeRebox(rebox, adaptor, destBoxTy, dest, base,2312 /*lbounds*/ {}, inputExtents, inputStrides,2313 rewriter);2314 2315 // The slice is of the form array(i:j:k)[%component]. Compute new extents2316 // and strides.2317 llvm::SmallVector<mlir::Value> slicedExtents;2318 llvm::SmallVector<mlir::Value> slicedStrides;2319 mlir::Value one = fir::genConstantIndex(loc, idxTy, rewriter, 1);2320 const bool sliceHasOrigins = !rebox.getShift().empty();2321 unsigned sliceOps = rebox.getSliceOperandIndex();2322 unsigned shiftOps = rebox.getShiftOperandIndex();2323 auto strideOps = inputStrides.begin();2324 const unsigned inputRank = inputStrides.size();2325 for (unsigned i = 0; i < inputRank;2326 ++i, ++strideOps, ++shiftOps, sliceOps += 3) {2327 mlir::Value sliceLb =2328 integerCast(loc, rewriter, idxTy, operands[sliceOps]);2329 mlir::Value inputStride = *strideOps; // already idxTy2330 // Apply origin shift: base += (lb-shift)*input_stride2331 mlir::Value sliceOrigin =2332 sliceHasOrigins2333 ? integerCast(loc, rewriter, idxTy, operands[shiftOps])2334 : one;2335 mlir::Value diff =2336 mlir::LLVM::SubOp::create(rewriter, loc, idxTy, sliceLb, sliceOrigin);2337 mlir::Value offset =2338 mlir::LLVM::MulOp::create(rewriter, loc, idxTy, diff, inputStride);2339 // Strides from the fir.box are in bytes.2340 base = genGEP(loc, byteTy, rewriter, base, offset);2341 // Apply upper bound and step if this is a triplet. Otherwise, the2342 // dimension is dropped and no extents/strides are computed.2343 mlir::Value upper = operands[sliceOps + 1];2344 const bool isTripletSlice =2345 !mlir::isa_and_nonnull<mlir::LLVM::UndefOp>(upper.getDefiningOp());2346 if (isTripletSlice) {2347 mlir::Value step =2348 integerCast(loc, rewriter, idxTy, operands[sliceOps + 2]);2349 // extent = ub-lb+step/step2350 mlir::Value sliceUb = integerCast(loc, rewriter, idxTy, upper);2351 mlir::Value extent = computeTripletExtent(rewriter, loc, sliceLb,2352 sliceUb, step, zero, idxTy);2353 slicedExtents.emplace_back(extent);2354 // stride = step*input_stride2355 mlir::Value stride =2356 mlir::LLVM::MulOp::create(rewriter, loc, idxTy, step, inputStride);2357 slicedStrides.emplace_back(stride);2358 }2359 }2360 return finalizeRebox(rebox, adaptor, destBoxTy, dest, base,2361 /*lbounds*/ {}, slicedExtents, slicedStrides,2362 rewriter);2363 }2364 2365 /// Apply a new shape to the data described by a box given the base address,2366 /// extents and strides of the box.2367 llvm::LogicalResult2368 reshapeBox(fir::cg::XReboxOp rebox, OpAdaptor adaptor, mlir::Type destBoxTy,2369 mlir::Value dest, mlir::Value base, mlir::ValueRange inputExtents,2370 mlir::ValueRange inputStrides, mlir::ValueRange operands,2371 mlir::ConversionPatternRewriter &rewriter) const {2372 mlir::ValueRange reboxShifts{2373 operands.begin() + rebox.getShiftOperandIndex(),2374 operands.begin() + rebox.getShiftOperandIndex() +2375 rebox.getShift().size()};2376 if (rebox.getShape().empty()) {2377 // Only setting new lower bounds.2378 return finalizeRebox(rebox, adaptor, destBoxTy, dest, base, reboxShifts,2379 inputExtents, inputStrides, rewriter);2380 }2381 2382 mlir::Location loc = rebox.getLoc();2383 2384 llvm::SmallVector<mlir::Value> newStrides;2385 llvm::SmallVector<mlir::Value> newExtents;2386 mlir::Type idxTy = lowerTy().indexType();2387 // First stride from input box is kept. The rest is assumed contiguous2388 // (it is not possible to reshape otherwise). If the input is scalar,2389 // which may be OK if all new extents are ones, the stride does not2390 // matter, use one.2391 mlir::Value stride = inputStrides.empty()2392 ? fir::genConstantIndex(loc, idxTy, rewriter, 1)2393 : inputStrides[0];2394 for (unsigned i = 0; i < rebox.getShape().size(); ++i) {2395 mlir::Value rawExtent = operands[rebox.getShapeOperandIndex() + i];2396 mlir::Value extent = integerCast(loc, rewriter, idxTy, rawExtent);2397 newExtents.emplace_back(extent);2398 newStrides.emplace_back(stride);2399 // nextStride = extent * stride;2400 stride = mlir::LLVM::MulOp::create(rewriter, loc, idxTy, extent, stride);2401 }2402 return finalizeRebox(rebox, adaptor, destBoxTy, dest, base, reboxShifts,2403 newExtents, newStrides, rewriter);2404 }2405 2406 /// Return scalar element type of the input box.2407 static mlir::Type getInputEleTy(fir::cg::XReboxOp rebox) {2408 auto ty = fir::dyn_cast_ptrOrBoxEleTy(rebox.getBox().getType());2409 if (auto seqTy = mlir::dyn_cast<fir::SequenceType>(ty))2410 return seqTy.getEleTy();2411 return ty;2412 }2413};2414 2415/// Lower `fir.emboxproc` operation. Creates a procedure box.2416/// TODO: Part of supporting Fortran 2003 procedure pointers.2417struct EmboxProcOpConversion : public fir::FIROpConversion<fir::EmboxProcOp> {2418 using FIROpConversion::FIROpConversion;2419 2420 llvm::LogicalResult2421 matchAndRewrite(fir::EmboxProcOp emboxproc, OpAdaptor adaptor,2422 mlir::ConversionPatternRewriter &rewriter) const override {2423 TODO(emboxproc.getLoc(), "fir.emboxproc codegen");2424 return mlir::failure();2425 }2426};2427 2428// Code shared between insert_value and extract_value Ops.2429struct ValueOpCommon {2430 // Translate the arguments pertaining to any multidimensional array to2431 // row-major order for LLVM-IR.2432 static void toRowMajor(llvm::SmallVectorImpl<int64_t> &indices,2433 mlir::Type ty) {2434 assert(ty && "type is null");2435 const auto end = indices.size();2436 for (std::remove_const_t<decltype(end)> i = 0; i < end; ++i) {2437 if (auto seq = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(ty)) {2438 const auto dim = getDimension(seq);2439 if (dim > 1) {2440 auto ub = std::min(i + dim, end);2441 std::reverse(indices.begin() + i, indices.begin() + ub);2442 i += dim - 1;2443 }2444 ty = getArrayElementType(seq);2445 } else if (auto st = mlir::dyn_cast<mlir::LLVM::LLVMStructType>(ty)) {2446 ty = st.getBody()[indices[i]];2447 } else {2448 llvm_unreachable("index into invalid type");2449 }2450 }2451 }2452 2453 static llvm::SmallVector<int64_t>2454 collectIndices(mlir::ConversionPatternRewriter &rewriter,2455 mlir::ArrayAttr arrAttr) {2456 llvm::SmallVector<int64_t> indices;2457 for (auto i = arrAttr.begin(), e = arrAttr.end(); i != e; ++i) {2458 if (auto intAttr = mlir::dyn_cast<mlir::IntegerAttr>(*i)) {2459 indices.push_back(intAttr.getInt());2460 } else {2461 auto fieldName = mlir::cast<mlir::StringAttr>(*i).getValue();2462 ++i;2463 auto ty = mlir::cast<mlir::TypeAttr>(*i).getValue();2464 auto index = mlir::cast<fir::RecordType>(ty).getFieldIndex(fieldName);2465 indices.push_back(index);2466 }2467 }2468 return indices;2469 }2470 2471private:2472 static mlir::Type getArrayElementType(mlir::LLVM::LLVMArrayType ty) {2473 auto eleTy = ty.getElementType();2474 while (auto arrTy = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(eleTy))2475 eleTy = arrTy.getElementType();2476 return eleTy;2477 }2478};2479 2480namespace {2481/// Extract a subobject value from an ssa-value of aggregate type2482struct ExtractValueOpConversion2483 : public fir::FIROpAndTypeConversion<fir::ExtractValueOp>,2484 public ValueOpCommon {2485 using FIROpAndTypeConversion::FIROpAndTypeConversion;2486 2487 llvm::LogicalResult2488 doRewrite(fir::ExtractValueOp extractVal, mlir::Type ty, OpAdaptor adaptor,2489 mlir::ConversionPatternRewriter &rewriter) const override {2490 mlir::ValueRange operands = adaptor.getOperands();2491 auto indices = collectIndices(rewriter, extractVal.getCoor());2492 toRowMajor(indices, operands[0].getType());2493 rewriter.replaceOpWithNewOp<mlir::LLVM::ExtractValueOp>(2494 extractVal, operands[0], indices);2495 return mlir::success();2496 }2497};2498 2499/// InsertValue is the generalized instruction for the composition of new2500/// aggregate type values.2501struct InsertValueOpConversion2502 : public mlir::OpConversionPattern<fir::InsertValueOp>,2503 public ValueOpCommon {2504 using OpConversionPattern::OpConversionPattern;2505 2506 llvm::LogicalResult2507 matchAndRewrite(fir::InsertValueOp insertVal, OpAdaptor adaptor,2508 mlir::ConversionPatternRewriter &rewriter) const override {2509 mlir::ValueRange operands = adaptor.getOperands();2510 auto indices = collectIndices(rewriter, insertVal.getCoor());2511 toRowMajor(indices, operands[0].getType());2512 rewriter.replaceOpWithNewOp<mlir::LLVM::InsertValueOp>(2513 insertVal, operands[0], operands[1], indices);2514 return mlir::success();2515 }2516};2517 2518/// InsertOnRange inserts a value into a sequence over a range of offsets.2519struct InsertOnRangeOpConversion2520 : public fir::FIROpAndTypeConversion<fir::InsertOnRangeOp> {2521 using FIROpAndTypeConversion::FIROpAndTypeConversion;2522 2523 // Increments an array of subscripts in a row major fasion.2524 void incrementSubscripts(llvm::ArrayRef<int64_t> dims,2525 llvm::SmallVectorImpl<int64_t> &subscripts) const {2526 for (size_t i = dims.size(); i > 0; --i) {2527 if (++subscripts[i - 1] < dims[i - 1]) {2528 return;2529 }2530 subscripts[i - 1] = 0;2531 }2532 }2533 2534 llvm::LogicalResult2535 doRewrite(fir::InsertOnRangeOp range, mlir::Type ty, OpAdaptor adaptor,2536 mlir::ConversionPatternRewriter &rewriter) const override {2537 2538 auto arrayType = adaptor.getSeq().getType();2539 2540 // Iteratively extract the array dimensions from the type.2541 llvm::SmallVector<std::int64_t> dims;2542 mlir::Type type = arrayType;2543 while (auto t = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(type)) {2544 dims.push_back(t.getNumElements());2545 type = t.getElementType();2546 }2547 2548 // Avoid generating long insert chain that are very slow to fold back2549 // (which is required in globals when later generating LLVM IR). Attempt to2550 // fold the inserted element value to an attribute and build an ArrayAttr2551 // for the resulting array.2552 if (range.isFullRange()) {2553 llvm::FailureOr<mlir::Attribute> cst =2554 fir::tryFoldingLLVMInsertChain(adaptor.getVal(), rewriter);2555 if (llvm::succeeded(cst)) {2556 mlir::Attribute dimVal = *cst;2557 for (auto dim : llvm::reverse(dims)) {2558 // Use std::vector in case the number of elements is big.2559 std::vector<mlir::Attribute> elements(dim, dimVal);2560 dimVal = mlir::ArrayAttr::get(range.getContext(), elements);2561 }2562 // Replace insert chain with constant.2563 rewriter.replaceOpWithNewOp<mlir::LLVM::ConstantOp>(range, arrayType,2564 dimVal);2565 return mlir::success();2566 }2567 }2568 2569 // The inserted value cannot be folded to an attribute, turn the2570 // insert_range into an llvm.insertvalue chain.2571 llvm::SmallVector<std::int64_t> lBounds;2572 llvm::SmallVector<std::int64_t> uBounds;2573 2574 // Unzip the upper and lower bound and convert to a row major format.2575 mlir::DenseIntElementsAttr coor = range.getCoor();2576 auto reversedCoor = llvm::reverse(coor.getValues<int64_t>());2577 for (auto i = reversedCoor.begin(), e = reversedCoor.end(); i != e; ++i) {2578 uBounds.push_back(*i++);2579 lBounds.push_back(*i);2580 }2581 2582 auto &subscripts = lBounds;2583 auto loc = range.getLoc();2584 mlir::Value lastOp = adaptor.getSeq();2585 mlir::Value insertVal = adaptor.getVal();2586 2587 while (subscripts != uBounds) {2588 lastOp = mlir::LLVM::InsertValueOp::create(rewriter, loc, lastOp,2589 insertVal, subscripts);2590 2591 incrementSubscripts(dims, subscripts);2592 }2593 2594 rewriter.replaceOpWithNewOp<mlir::LLVM::InsertValueOp>(2595 range, lastOp, insertVal, subscripts);2596 2597 return mlir::success();2598 }2599};2600} // namespace2601 2602namespace {2603/// XArrayCoor is the address arithmetic on a dynamically shaped, sliced,2604/// shifted etc. array.2605/// (See the static restriction on coordinate_of.) array_coor determines the2606/// coordinate (location) of a specific element.2607struct XArrayCoorOpConversion2608 : public fir::FIROpAndTypeConversion<fir::cg::XArrayCoorOp> {2609 using FIROpAndTypeConversion::FIROpAndTypeConversion;2610 2611 llvm::LogicalResult2612 doRewrite(fir::cg::XArrayCoorOp coor, mlir::Type llvmPtrTy, OpAdaptor adaptor,2613 mlir::ConversionPatternRewriter &rewriter) const override {2614 auto loc = coor.getLoc();2615 mlir::ValueRange operands = adaptor.getOperands();2616 unsigned rank = coor.getRank();2617 assert(coor.getIndices().size() == rank);2618 assert(coor.getShape().empty() || coor.getShape().size() == rank);2619 assert(coor.getShift().empty() || coor.getShift().size() == rank);2620 assert(coor.getSlice().empty() || coor.getSlice().size() == 3 * rank);2621 mlir::Type idxTy = lowerTy().indexType();2622 unsigned indexOffset = coor.getIndicesOperandIndex();2623 unsigned shapeOffset = coor.getShapeOperandIndex();2624 unsigned shiftOffset = coor.getShiftOperandIndex();2625 unsigned sliceOffset = coor.getSliceOperandIndex();2626 auto sliceOps = coor.getSlice().begin();2627 mlir::Value one = fir::genConstantIndex(loc, idxTy, rewriter, 1);2628 mlir::Value prevExt = one;2629 mlir::Value offset = fir::genConstantIndex(loc, idxTy, rewriter, 0);2630 const bool isShifted = !coor.getShift().empty();2631 const bool isSliced = !coor.getSlice().empty();2632 const bool baseIsBoxed =2633 mlir::isa<fir::BaseBoxType>(coor.getMemref().getType());2634 TypePair baseBoxTyPair =2635 baseIsBoxed ? getBoxTypePair(coor.getMemref().getType()) : TypePair{};2636 mlir::LLVM::IntegerOverflowFlags nsw =2637 mlir::LLVM::IntegerOverflowFlags::nsw;2638 2639 // For each dimension of the array, generate the offset calculation.2640 for (unsigned i = 0; i < rank; ++i, ++indexOffset, ++shapeOffset,2641 ++shiftOffset, sliceOffset += 3, sliceOps += 3) {2642 mlir::Value index =2643 integerCast(loc, rewriter, idxTy, operands[indexOffset]);2644 mlir::Value lb =2645 isShifted ? integerCast(loc, rewriter, idxTy, operands[shiftOffset])2646 : one;2647 mlir::Value step = one;2648 bool normalSlice = isSliced;2649 // Compute zero based index in dimension i of the element, applying2650 // potential triplets and lower bounds.2651 if (isSliced) {2652 mlir::Value originalUb = *(sliceOps + 1);2653 normalSlice =2654 !mlir::isa_and_nonnull<fir::UndefOp>(originalUb.getDefiningOp());2655 if (normalSlice)2656 step = integerCast(loc, rewriter, idxTy, operands[sliceOffset + 2]);2657 }2658 auto idx =2659 mlir::LLVM::SubOp::create(rewriter, loc, idxTy, index, lb, nsw);2660 mlir::Value diff =2661 mlir::LLVM::MulOp::create(rewriter, loc, idxTy, idx, step, nsw);2662 if (normalSlice) {2663 mlir::Value sliceLb =2664 integerCast(loc, rewriter, idxTy, operands[sliceOffset]);2665 auto adj =2666 mlir::LLVM::SubOp::create(rewriter, loc, idxTy, sliceLb, lb, nsw);2667 diff = mlir::LLVM::AddOp::create(rewriter, loc, idxTy, diff, adj, nsw);2668 }2669 // Update the offset given the stride and the zero based index `diff`2670 // that was just computed.2671 if (baseIsBoxed) {2672 // Use stride in bytes from the descriptor.2673 mlir::Value stride =2674 getStrideFromBox(loc, baseBoxTyPair, operands[0], i, rewriter);2675 auto sc =2676 mlir::LLVM::MulOp::create(rewriter, loc, idxTy, diff, stride, nsw);2677 offset =2678 mlir::LLVM::AddOp::create(rewriter, loc, idxTy, sc, offset, nsw);2679 } else {2680 // Use stride computed at last iteration.2681 auto sc =2682 mlir::LLVM::MulOp::create(rewriter, loc, idxTy, diff, prevExt, nsw);2683 offset =2684 mlir::LLVM::AddOp::create(rewriter, loc, idxTy, sc, offset, nsw);2685 // Compute next stride assuming contiguity of the base array2686 // (in element number).2687 auto nextExt = integerCast(loc, rewriter, idxTy, operands[shapeOffset]);2688 prevExt = mlir::LLVM::MulOp::create(rewriter, loc, idxTy, prevExt,2689 nextExt, nsw);2690 }2691 }2692 2693 // Add computed offset to the base address.2694 if (baseIsBoxed) {2695 // Working with byte offsets. The base address is read from the fir.box.2696 // and used in i8* GEP to do the pointer arithmetic.2697 mlir::Type byteTy = ::getI8Type(coor.getContext());2698 mlir::Value base =2699 getBaseAddrFromBox(loc, baseBoxTyPair, operands[0], rewriter);2700 llvm::SmallVector<mlir::LLVM::GEPArg> args{offset};2701 auto addr = mlir::LLVM::GEPOp::create(rewriter, loc, llvmPtrTy, byteTy,2702 base, args);2703 if (coor.getSubcomponent().empty()) {2704 rewriter.replaceOp(coor, addr);2705 return mlir::success();2706 }2707 // Cast the element address from void* to the derived type so that the2708 // derived type members can be addresses via a GEP using the index of2709 // components.2710 mlir::Type elementType =2711 getLlvmObjectTypeFromBoxType(coor.getMemref().getType());2712 while (auto arrayTy =2713 mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(elementType))2714 elementType = arrayTy.getElementType();2715 args.clear();2716 args.push_back(0);2717 if (!coor.getLenParams().empty()) {2718 // If type parameters are present, then we don't want to use a GEPOp2719 // as below, as the LLVM struct type cannot be statically defined.2720 TODO(loc, "derived type with type parameters");2721 }2722 llvm::SmallVector<mlir::Value> indices = convertSubcomponentIndices(2723 loc, elementType,2724 operands.slice(coor.getSubcomponentOperandIndex(),2725 coor.getSubcomponent().size()));2726 args.append(indices.begin(), indices.end());2727 rewriter.replaceOpWithNewOp<mlir::LLVM::GEPOp>(coor, llvmPtrTy,2728 elementType, addr, args);2729 return mlir::success();2730 }2731 2732 // The array was not boxed, so it must be contiguous. offset is therefore an2733 // element offset and the base type is kept in the GEP unless the element2734 // type size is itself dynamic.2735 mlir::Type objectTy = fir::unwrapRefType(coor.getMemref().getType());2736 mlir::Type eleType = fir::unwrapSequenceType(objectTy);2737 mlir::Type gepObjectType = convertType(eleType);2738 llvm::SmallVector<mlir::LLVM::GEPArg> args;2739 if (coor.getSubcomponent().empty()) {2740 // No subcomponent.2741 if (!coor.getLenParams().empty()) {2742 // Type parameters. Adjust element size explicitly.2743 auto eleTy = fir::dyn_cast_ptrEleTy(coor.getType());2744 assert(eleTy && "result must be a reference-like type");2745 if (fir::characterWithDynamicLen(eleTy)) {2746 assert(coor.getLenParams().size() == 1);2747 auto length = integerCast(loc, rewriter, idxTy,2748 operands[coor.getLenParamsOperandIndex()]);2749 offset = mlir::LLVM::MulOp::create(rewriter, loc, idxTy, offset,2750 length, nsw);2751 } else {2752 TODO(loc, "compute size of derived type with type parameters");2753 }2754 }2755 args.push_back(offset);2756 } else {2757 // There are subcomponents.2758 args.push_back(offset);2759 llvm::SmallVector<mlir::Value> indices = convertSubcomponentIndices(2760 loc, gepObjectType,2761 operands.slice(coor.getSubcomponentOperandIndex(),2762 coor.getSubcomponent().size()));2763 args.append(indices.begin(), indices.end());2764 }2765 rewriter.replaceOpWithNewOp<mlir::LLVM::GEPOp>(2766 coor, llvmPtrTy, gepObjectType, adaptor.getMemref(), args);2767 return mlir::success();2768 }2769};2770} // namespace2771 2772/// Convert to (memory) reference to a reference to a subobject.2773/// The coordinate_of op is a Swiss army knife operation that can be used on2774/// (memory) references to records, arrays, complex, etc. as well as boxes.2775/// With unboxed arrays, there is the restriction that the array have a static2776/// shape in all but the last column.2777struct CoordinateOpConversion2778 : public fir::FIROpAndTypeConversion<fir::CoordinateOp> {2779 using FIROpAndTypeConversion::FIROpAndTypeConversion;2780 2781 llvm::LogicalResult2782 doRewrite(fir::CoordinateOp coor, mlir::Type ty, OpAdaptor adaptor,2783 mlir::ConversionPatternRewriter &rewriter) const override {2784 mlir::ValueRange operands = adaptor.getOperands();2785 2786 mlir::Location loc = coor.getLoc();2787 mlir::Value base = operands[0];2788 mlir::Type baseObjectTy = coor.getBaseType();2789 mlir::Type objectTy = fir::dyn_cast_ptrOrBoxEleTy(baseObjectTy);2790 assert(objectTy && "fir.coordinate_of expects a reference type");2791 mlir::Type llvmObjectTy = convertType(objectTy);2792 2793 // Complex type - basically, extract the real or imaginary part2794 // FIXME: double check why this is done before the fir.box case below.2795 if (fir::isa_complex(objectTy)) {2796 mlir::Value gep =2797 genGEP(loc, llvmObjectTy, rewriter, base, 0, operands[1]);2798 rewriter.replaceOp(coor, gep);2799 return mlir::success();2800 }2801 2802 // Boxed type - get the base pointer from the box2803 if (mlir::dyn_cast<fir::BaseBoxType>(baseObjectTy))2804 return doRewriteBox(coor, operands, loc, rewriter);2805 2806 // Reference, pointer or a heap type2807 if (mlir::isa<fir::ReferenceType, fir::PointerType, fir::HeapType>(2808 baseObjectTy))2809 return doRewriteRefOrPtr(coor, llvmObjectTy, operands, loc, rewriter);2810 2811 return rewriter.notifyMatchFailure(2812 coor, "fir.coordinate_of base operand has unsupported type");2813 }2814 2815 static unsigned getFieldNumber(fir::RecordType ty, mlir::Value op) {2816 return fir::hasDynamicSize(ty)2817 ? op.getDefiningOp()2818 ->getAttrOfType<mlir::IntegerAttr>("field")2819 .getInt()2820 : getConstantIntValue(op);2821 }2822 2823 static bool hasSubDimensions(mlir::Type type) {2824 return mlir::isa<fir::SequenceType, fir::RecordType, mlir::TupleType>(type);2825 }2826 2827 // Helper structure to analyze the CoordinateOp path and decide if and how2828 // the GEP should be generated for it.2829 struct ShapeAnalysis {2830 bool hasKnownShape;2831 bool columnIsDeferred;2832 };2833 2834 /// Walk the abstract memory layout and determine if the path traverses any2835 /// array types with unknown shape. Return true iff all the array types have a2836 /// constant shape along the path.2837 /// TODO: move the verification logic into the verifier.2838 static std::optional<ShapeAnalysis>2839 arraysHaveKnownShape(mlir::Type type, fir::CoordinateOp coor) {2840 fir::CoordinateIndicesAdaptor indices = coor.getIndices();2841 auto begin = indices.begin();2842 bool hasKnownShape = true;2843 bool columnIsDeferred = false;2844 for (auto it = begin, end = indices.end(); it != end;) {2845 if (auto arrTy = mlir::dyn_cast<fir::SequenceType>(type)) {2846 bool addressingStart = (it == begin);2847 unsigned arrayDim = arrTy.getDimension();2848 for (auto dimExtent : llvm::enumerate(arrTy.getShape())) {2849 if (dimExtent.value() == fir::SequenceType::getUnknownExtent()) {2850 hasKnownShape = false;2851 if (addressingStart && dimExtent.index() + 1 == arrayDim) {2852 // If this point was reached, the raws of the first array have2853 // constant extents.2854 columnIsDeferred = true;2855 } else {2856 // One of the array dimension that is not the column of the first2857 // array has dynamic extent. It will not possible to do2858 // code generation for the CoordinateOp if the base is not a2859 // fir.box containing the value of that extent.2860 return ShapeAnalysis{false, false};2861 }2862 }2863 // There may be less operands than the array size if the2864 // fir.coordinate_of result is not an element but a sub-array.2865 if (it != end)2866 ++it;2867 }2868 type = arrTy.getEleTy();2869 continue;2870 }2871 if (auto strTy = mlir::dyn_cast<fir::RecordType>(type)) {2872 auto intAttr = llvm::dyn_cast<mlir::IntegerAttr>(*it);2873 if (!intAttr) {2874 mlir::emitError(coor.getLoc(),2875 "expected field name in fir.coordinate_of");2876 return std::nullopt;2877 }2878 type = strTy.getType(intAttr.getInt());2879 } else if (auto strTy = mlir::dyn_cast<mlir::TupleType>(type)) {2880 auto value = llvm::dyn_cast<mlir::Value>(*it);2881 if (!value) {2882 mlir::emitError(2883 coor.getLoc(),2884 "expected constant value to address tuple in fir.coordinate_of");2885 return std::nullopt;2886 }2887 type = strTy.getType(getConstantIntValue(value));2888 } else if (auto charType = mlir::dyn_cast<fir::CharacterType>(type)) {2889 // Addressing character in string. Fortran strings degenerate to arrays2890 // in LLVM, so they are handled like arrays of characters here.2891 if (charType.getLen() == fir::CharacterType::unknownLen())2892 return ShapeAnalysis{false, true};2893 type = fir::CharacterType::getSingleton(charType.getContext(),2894 charType.getFKind());2895 }2896 ++it;2897 }2898 return ShapeAnalysis{hasKnownShape, columnIsDeferred};2899 }2900 2901private:2902 llvm::LogicalResult2903 doRewriteBox(fir::CoordinateOp coor, mlir::ValueRange operands,2904 mlir::Location loc,2905 mlir::ConversionPatternRewriter &rewriter) const {2906 mlir::Type boxObjTy = coor.getBaseType();2907 assert(mlir::dyn_cast<fir::BaseBoxType>(boxObjTy) &&2908 "This is not a `fir.box`");2909 TypePair boxTyPair = getBoxTypePair(boxObjTy);2910 2911 mlir::Value boxBaseAddr = operands[0];2912 2913 // 1. SPECIAL CASE (uses `fir.len_param_index`):2914 // %box = ... : !fir.box<!fir.type<derived{len1:i32}>>2915 // %lenp = fir.len_param_index len1, !fir.type<derived{len1:i32}>2916 // %addr = coordinate_of %box, %lenp2917 if (coor.getNumOperands() == 2) {2918 mlir::Operation *coordinateDef =2919 (*coor.getCoor().begin()).getDefiningOp();2920 if (mlir::isa_and_nonnull<fir::LenParamIndexOp>(coordinateDef))2921 TODO(loc,2922 "fir.coordinate_of - fir.len_param_index is not supported yet");2923 }2924 2925 // 2. GENERAL CASE:2926 // 2.1. (`fir.array`)2927 // %box = ... : !fix.box<!fir.array<?xU>>2928 // %idx = ... : index2929 // %resultAddr = coordinate_of %box, %idx : !fir.ref<U>2930 // 2.2 (`fir.derived`)2931 // %box = ... : !fix.box<!fir.type<derived_type{field_1:i32}>>2932 // %idx = ... : i322933 // %resultAddr = coordinate_of %box, %idx : !fir.ref<i32>2934 // 2.3 (`fir.derived` inside `fir.array`)2935 // %box = ... : !fir.box<!fir.array<10 x !fir.type<derived_1{field_1:f32,2936 // field_2:f32}>>> %idx1 = ... : index %idx2 = ... : i32 %resultAddr =2937 // coordinate_of %box, %idx1, %idx2 : !fir.ref<f32>2938 // 2.4. TODO: Either document or disable any other case that the following2939 // implementation might convert.2940 mlir::Value resultAddr =2941 getBaseAddrFromBox(loc, boxTyPair, boxBaseAddr, rewriter);2942 // Component Type2943 auto cpnTy = fir::dyn_cast_ptrOrBoxEleTy(boxObjTy);2944 mlir::Type llvmPtrTy = ::getLlvmPtrType(coor.getContext());2945 mlir::Type byteTy = ::getI8Type(coor.getContext());2946 mlir::LLVM::IntegerOverflowFlags nsw =2947 mlir::LLVM::IntegerOverflowFlags::nsw;2948 2949 int nextIndexValue = 1;2950 fir::CoordinateIndicesAdaptor indices = coor.getIndices();2951 for (auto it = indices.begin(), end = indices.end(); it != end;) {2952 if (auto arrTy = mlir::dyn_cast<fir::SequenceType>(cpnTy)) {2953 if (it != indices.begin())2954 TODO(loc, "fir.array nested inside other array and/or derived type");2955 // Applies byte strides from the box. Ignore lower bound from box2956 // since fir.coordinate_of indexes are zero based. Lowering takes care2957 // of lower bound aspects. This both accounts for dynamically sized2958 // types and non contiguous arrays.2959 auto idxTy = lowerTy().indexType();2960 mlir::Value off = fir::genConstantIndex(loc, idxTy, rewriter, 0);2961 unsigned arrayDim = arrTy.getDimension();2962 for (unsigned dim = 0; dim < arrayDim && it != end; ++dim, ++it) {2963 mlir::Value stride =2964 getStrideFromBox(loc, boxTyPair, operands[0], dim, rewriter);2965 auto sc = mlir::LLVM::MulOp::create(rewriter, loc, idxTy,2966 operands[nextIndexValue + dim],2967 stride, nsw);2968 off = mlir::LLVM::AddOp::create(rewriter, loc, idxTy, sc, off, nsw);2969 }2970 nextIndexValue += arrayDim;2971 resultAddr = mlir::LLVM::GEPOp::create(2972 rewriter, loc, llvmPtrTy, byteTy, resultAddr,2973 llvm::ArrayRef<mlir::LLVM::GEPArg>{off});2974 cpnTy = arrTy.getEleTy();2975 } else if (auto recTy = mlir::dyn_cast<fir::RecordType>(cpnTy)) {2976 auto intAttr = llvm::dyn_cast<mlir::IntegerAttr>(*it);2977 if (!intAttr)2978 return mlir::emitError(loc,2979 "expected field name in fir.coordinate_of");2980 int fieldIndex = intAttr.getInt();2981 ++it;2982 cpnTy = recTy.getType(fieldIndex);2983 auto llvmRecTy = lowerTy().convertType(recTy);2984 resultAddr = mlir::LLVM::GEPOp::create(2985 rewriter, loc, llvmPtrTy, llvmRecTy, resultAddr,2986 llvm::ArrayRef<mlir::LLVM::GEPArg>{0, fieldIndex});2987 } else {2988 fir::emitFatalError(loc, "unexpected type in coordinate_of");2989 }2990 }2991 2992 rewriter.replaceOp(coor, resultAddr);2993 return mlir::success();2994 }2995 2996 llvm::LogicalResult2997 doRewriteRefOrPtr(fir::CoordinateOp coor, mlir::Type llvmObjectTy,2998 mlir::ValueRange operands, mlir::Location loc,2999 mlir::ConversionPatternRewriter &rewriter) const {3000 mlir::Type baseObjectTy = coor.getBaseType();3001 3002 // Component Type3003 mlir::Type cpnTy = fir::dyn_cast_ptrOrBoxEleTy(baseObjectTy);3004 3005 const std::optional<ShapeAnalysis> shapeAnalysis =3006 arraysHaveKnownShape(cpnTy, coor);3007 if (!shapeAnalysis)3008 return mlir::failure();3009 3010 if (fir::hasDynamicSize(fir::unwrapSequenceType(cpnTy)))3011 return mlir::emitError(3012 loc, "fir.coordinate_of with a dynamic element size is unsupported");3013 3014 if (shapeAnalysis->hasKnownShape || shapeAnalysis->columnIsDeferred) {3015 llvm::SmallVector<mlir::LLVM::GEPArg> offs;3016 if (shapeAnalysis->hasKnownShape) {3017 offs.push_back(0);3018 }3019 // Else, only the column is `?` and we can simply place the column value3020 // in the 0-th GEP position.3021 3022 std::optional<int> dims;3023 llvm::SmallVector<mlir::Value> arrIdx;3024 int nextIndexValue = 1;3025 for (auto index : coor.getIndices()) {3026 if (auto intAttr = llvm::dyn_cast<mlir::IntegerAttr>(index)) {3027 // Addressing derived type component.3028 auto recordType = llvm::dyn_cast<fir::RecordType>(cpnTy);3029 if (!recordType)3030 return mlir::emitError(3031 loc,3032 "fir.coordinate base type is not consistent with operands");3033 int fieldId = intAttr.getInt();3034 cpnTy = recordType.getType(fieldId);3035 offs.push_back(fieldId);3036 continue;3037 }3038 // Value index (addressing array, tuple, or complex part).3039 mlir::Value indexValue = operands[nextIndexValue++];3040 if (auto tupTy = mlir::dyn_cast<mlir::TupleType>(cpnTy)) {3041 cpnTy = tupTy.getType(getConstantIntValue(indexValue));3042 offs.push_back(indexValue);3043 } else {3044 if (!dims) {3045 if (auto arrayType = llvm::dyn_cast<fir::SequenceType>(cpnTy)) {3046 // Starting addressing array or array component.3047 dims = arrayType.getDimension();3048 cpnTy = arrayType.getElementType();3049 }3050 }3051 if (dims) {3052 arrIdx.push_back(indexValue);3053 if (--(*dims) == 0) {3054 // Append array range in reverse (FIR arrays are column-major).3055 offs.append(arrIdx.rbegin(), arrIdx.rend());3056 arrIdx.clear();3057 dims.reset();3058 }3059 } else {3060 offs.push_back(indexValue);3061 }3062 }3063 }3064 // It is possible the fir.coordinate_of result is a sub-array, in which3065 // case there may be some "unfinished" array indices to reverse and push.3066 if (!arrIdx.empty())3067 offs.append(arrIdx.rbegin(), arrIdx.rend());3068 3069 mlir::Value base = operands[0];3070 mlir::Value retval = genGEP(loc, llvmObjectTy, rewriter, base, offs);3071 rewriter.replaceOp(coor, retval);3072 return mlir::success();3073 }3074 3075 return mlir::emitError(3076 loc, "fir.coordinate_of base operand has unsupported type");3077 }3078};3079 3080/// Convert `fir.field_index`. The conversion depends on whether the size of3081/// the record is static or dynamic.3082struct FieldIndexOpConversion : public fir::FIROpConversion<fir::FieldIndexOp> {3083 using FIROpConversion::FIROpConversion;3084 3085 // NB: most field references should be resolved by this point3086 llvm::LogicalResult3087 matchAndRewrite(fir::FieldIndexOp field, OpAdaptor adaptor,3088 mlir::ConversionPatternRewriter &rewriter) const override {3089 auto recTy = mlir::cast<fir::RecordType>(field.getOnType());3090 unsigned index = recTy.getFieldIndex(field.getFieldId());3091 3092 if (!fir::hasDynamicSize(recTy)) {3093 // Derived type has compile-time constant layout. Return index of the3094 // component type in the parent type (to be used in GEP).3095 rewriter.replaceOp(field, mlir::ValueRange{genConstantOffset(3096 field.getLoc(), rewriter, index)});3097 return mlir::success();3098 }3099 3100 // Derived type has compile-time constant layout. Call the compiler3101 // generated function to determine the byte offset of the field at runtime.3102 // This returns a non-constant.3103 mlir::FlatSymbolRefAttr symAttr = mlir::SymbolRefAttr::get(3104 field.getContext(), getOffsetMethodName(recTy, field.getFieldId()));3105 mlir::NamedAttribute callAttr = rewriter.getNamedAttr("callee", symAttr);3106 mlir::NamedAttribute fieldAttr = rewriter.getNamedAttr(3107 "field", mlir::IntegerAttr::get(lowerTy().indexType(), index));3108 rewriter.replaceOpWithNewOp<mlir::LLVM::CallOp>(3109 field, lowerTy().offsetType(), adaptor.getOperands(),3110 addLLVMOpBundleAttrs(rewriter, {callAttr, fieldAttr},3111 adaptor.getOperands().size()));3112 return mlir::success();3113 }3114 3115 // Re-Construct the name of the compiler generated method that calculates the3116 // offset3117 inline static std::string getOffsetMethodName(fir::RecordType recTy,3118 llvm::StringRef field) {3119 return recTy.getName().str() + "P." + field.str() + ".offset";3120 }3121};3122 3123/// Convert `fir.end`3124struct FirEndOpConversion : public fir::FIROpConversion<fir::FirEndOp> {3125 using FIROpConversion::FIROpConversion;3126 3127 llvm::LogicalResult3128 matchAndRewrite(fir::FirEndOp firEnd, OpAdaptor,3129 mlir::ConversionPatternRewriter &rewriter) const override {3130 TODO(firEnd.getLoc(), "fir.end codegen");3131 return mlir::failure();3132 }3133};3134 3135/// Lower `fir.type_desc` to a global addr.3136struct TypeDescOpConversion : public fir::FIROpConversion<fir::TypeDescOp> {3137 using FIROpConversion::FIROpConversion;3138 3139 llvm::LogicalResult3140 matchAndRewrite(fir::TypeDescOp typeDescOp, OpAdaptor adaptor,3141 mlir::ConversionPatternRewriter &rewriter) const override {3142 mlir::Type inTy = typeDescOp.getInType();3143 assert(mlir::isa<fir::RecordType>(inTy) && "expecting fir.type");3144 auto recordType = mlir::dyn_cast<fir::RecordType>(inTy);3145 auto module = typeDescOp.getOperation()->getParentOfType<mlir::ModuleOp>();3146 mlir::Value typeDesc = getTypeDescriptor(3147 module, rewriter, typeDescOp.getLoc(), recordType, this->options);3148 rewriter.replaceOp(typeDescOp, typeDesc);3149 return mlir::success();3150 }3151};3152 3153/// Lower `fir.has_value` operation to `llvm.return` operation.3154struct HasValueOpConversion3155 : public mlir::OpConversionPattern<fir::HasValueOp> {3156 using OpConversionPattern::OpConversionPattern;3157 3158 llvm::LogicalResult3159 matchAndRewrite(fir::HasValueOp op, OpAdaptor adaptor,3160 mlir::ConversionPatternRewriter &rewriter) const override {3161 rewriter.replaceOpWithNewOp<mlir::LLVM::ReturnOp>(op,3162 adaptor.getOperands());3163 return mlir::success();3164 }3165};3166 3167#ifndef NDEBUG3168// Check if attr's type is compatible with ty.3169//3170// This is done by comparing attr's element type, converted to LLVM type,3171// with ty's element type.3172//3173// Only integer and floating point (including complex) attributes are3174// supported. Also, attr is expected to have a TensorType and ty is expected3175// to be of LLVMArrayType. If any of the previous conditions is false, then3176// the specified attr and ty are not supported by this function and are3177// assumed to be compatible.3178static inline bool attributeTypeIsCompatible(mlir::MLIRContext *ctx,3179 mlir::Attribute attr,3180 mlir::Type ty) {3181 // Get attr's LLVM element type.3182 if (!attr)3183 return true;3184 auto intOrFpEleAttr = mlir::dyn_cast<mlir::DenseIntOrFPElementsAttr>(attr);3185 if (!intOrFpEleAttr)3186 return true;3187 auto tensorTy = mlir::dyn_cast<mlir::TensorType>(intOrFpEleAttr.getType());3188 if (!tensorTy)3189 return true;3190 mlir::Type attrEleTy =3191 mlir::LLVMTypeConverter(ctx).convertType(tensorTy.getElementType());3192 3193 // Get ty's element type.3194 auto arrTy = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(ty);3195 if (!arrTy)3196 return true;3197 mlir::Type eleTy = arrTy.getElementType();3198 while ((arrTy = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(eleTy)))3199 eleTy = arrTy.getElementType();3200 3201 return attrEleTy == eleTy;3202}3203#endif3204 3205/// Lower `fir.global` operation to `llvm.global` operation.3206/// `fir.insert_on_range` operations are replaced with constant dense attribute3207/// if they are applied on the full range.3208struct GlobalOpConversion : public fir::FIROpConversion<fir::GlobalOp> {3209 using FIROpConversion::FIROpConversion;3210 3211 llvm::LogicalResult3212 matchAndRewrite(fir::GlobalOp global, OpAdaptor adaptor,3213 mlir::ConversionPatternRewriter &rewriter) const override {3214 3215 llvm::SmallVector<mlir::Attribute> dbgExprs;3216 3217 if (auto fusedLoc = mlir::dyn_cast<mlir::FusedLoc>(global.getLoc())) {3218 if (auto gvExprAttr = mlir::dyn_cast_if_present<mlir::ArrayAttr>(3219 fusedLoc.getMetadata())) {3220 for (auto attr : gvExprAttr.getAsRange<mlir::Attribute>())3221 if (auto dbgAttr =3222 mlir::dyn_cast<mlir::LLVM::DIGlobalVariableExpressionAttr>(3223 attr))3224 dbgExprs.push_back(dbgAttr);3225 }3226 }3227 3228 auto tyAttr = convertType(global.getType());3229 if (auto boxType = mlir::dyn_cast<fir::BaseBoxType>(global.getType()))3230 tyAttr = this->lowerTy().convertBoxTypeAsStruct(boxType);3231 auto loc = global.getLoc();3232 mlir::Attribute initAttr = global.getInitVal().value_or(mlir::Attribute());3233 assert(attributeTypeIsCompatible(global.getContext(), initAttr, tyAttr));3234 auto linkage = convertLinkage(global.getLinkName());3235 auto isConst = global.getConstant().has_value();3236 mlir::SymbolRefAttr comdat;3237 llvm::ArrayRef<mlir::NamedAttribute> attrs;3238 auto g = mlir::LLVM::GlobalOp::create(3239 rewriter, loc, tyAttr, isConst, linkage, global.getSymName(), initAttr,3240 0, getGlobalAddressSpace(rewriter), false, false, comdat, attrs,3241 dbgExprs);3242 3243 if (global.getAlignment() && *global.getAlignment() > 0)3244 g.setAlignment(*global.getAlignment());3245 3246 auto module = global->getParentOfType<mlir::ModuleOp>();3247 auto gpuMod = global->getParentOfType<mlir::gpu::GPUModuleOp>();3248 // Add comdat if necessary3249 if (fir::getTargetTriple(module).supportsCOMDAT() &&3250 (linkage == mlir::LLVM::Linkage::Linkonce ||3251 linkage == mlir::LLVM::Linkage::LinkonceODR) &&3252 !gpuMod) {3253 addComdat(g, rewriter, module);3254 }3255 3256 // Apply all non-Fir::GlobalOp attributes to the LLVM::GlobalOp, preserving3257 // them; whilst taking care not to apply attributes that are lowered in3258 // other ways.3259 llvm::SmallDenseSet<llvm::StringRef> elidedAttrsSet(3260 global.getAttributeNames().begin(), global.getAttributeNames().end());3261 for (auto &attr : global->getAttrs())3262 if (!elidedAttrsSet.contains(attr.getName().strref()))3263 g->setAttr(attr.getName(), attr.getValue());3264 3265 auto &gr = g.getInitializerRegion();3266 rewriter.inlineRegionBefore(global.getRegion(), gr, gr.end());3267 if (!gr.empty()) {3268 // Replace insert_on_range with a constant dense attribute if the3269 // initialization is on the full range.3270 auto insertOnRangeOps = gr.front().getOps<fir::InsertOnRangeOp>();3271 for (auto insertOp : insertOnRangeOps) {3272 if (insertOp.isFullRange()) {3273 auto seqTyAttr = convertType(insertOp.getType());3274 auto *op = insertOp.getVal().getDefiningOp();3275 auto constant = mlir::dyn_cast<mlir::arith::ConstantOp>(op);3276 if (!constant) {3277 auto convertOp = mlir::dyn_cast<fir::ConvertOp>(op);3278 if (!convertOp)3279 continue;3280 constant = mlir::cast<mlir::arith::ConstantOp>(3281 convertOp.getValue().getDefiningOp());3282 }3283 mlir::Type vecType = mlir::VectorType::get(3284 insertOp.getType().getShape(), constant.getType());3285 auto denseAttr = mlir::DenseElementsAttr::get(3286 mlir::cast<mlir::ShapedType>(vecType), constant.getValue());3287 rewriter.setInsertionPointAfter(insertOp);3288 rewriter.replaceOpWithNewOp<mlir::arith::ConstantOp>(3289 insertOp, seqTyAttr, denseAttr);3290 }3291 }3292 }3293 3294 if (global.getDataAttr() &&3295 *global.getDataAttr() == cuf::DataAttribute::Shared)3296 g.setAddrSpace(3297 static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Shared));3298 3299 if (global.getDataAttr() &&3300 *global.getDataAttr() == cuf::DataAttribute::Constant)3301 g.setAddrSpace(3302 static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Constant));3303 3304 rewriter.eraseOp(global);3305 return mlir::success();3306 }3307 3308 // TODO: String comparisons should be avoided. Replace linkName with an3309 // enumeration.3310 mlir::LLVM::Linkage3311 convertLinkage(std::optional<llvm::StringRef> optLinkage) const {3312 if (optLinkage) {3313 auto name = *optLinkage;3314 if (name == "internal")3315 return mlir::LLVM::Linkage::Internal;3316 if (name == "linkonce")3317 return mlir::LLVM::Linkage::Linkonce;3318 if (name == "linkonce_odr")3319 return mlir::LLVM::Linkage::LinkonceODR;3320 if (name == "common")3321 return mlir::LLVM::Linkage::Common;3322 if (name == "weak")3323 return mlir::LLVM::Linkage::Weak;3324 }3325 return mlir::LLVM::Linkage::External;3326 }3327 3328private:3329 static void addComdat(mlir::LLVM::GlobalOp &global,3330 mlir::ConversionPatternRewriter &rewriter,3331 mlir::ModuleOp module) {3332 const char *comdatName = "__llvm_comdat";3333 mlir::LLVM::ComdatOp comdatOp =3334 module.lookupSymbol<mlir::LLVM::ComdatOp>(comdatName);3335 if (!comdatOp) {3336 comdatOp =3337 mlir::LLVM::ComdatOp::create(rewriter, module.getLoc(), comdatName);3338 }3339 if (auto select = comdatOp.lookupSymbol<mlir::LLVM::ComdatSelectorOp>(3340 global.getSymName()))3341 return;3342 mlir::OpBuilder::InsertionGuard guard(rewriter);3343 rewriter.setInsertionPointToEnd(&comdatOp.getBody().back());3344 auto selectorOp = mlir::LLVM::ComdatSelectorOp::create(3345 rewriter, comdatOp.getLoc(), global.getSymName(),3346 mlir::LLVM::comdat::Comdat::Any);3347 global.setComdatAttr(mlir::SymbolRefAttr::get(3348 rewriter.getContext(), comdatName,3349 mlir::FlatSymbolRefAttr::get(selectorOp.getSymNameAttr())));3350 }3351};3352 3353/// `fir.load` --> `llvm.load`3354struct LoadOpConversion : public fir::FIROpConversion<fir::LoadOp> {3355 using FIROpConversion::FIROpConversion;3356 3357 llvm::LogicalResult3358 matchAndRewrite(fir::LoadOp load, OpAdaptor adaptor,3359 mlir::ConversionPatternRewriter &rewriter) const override {3360 3361 mlir::Type llvmLoadTy = convertObjectType(load.getType());3362 const bool isVolatile = fir::isa_volatile_type(load.getMemref().getType());3363 if (auto boxTy = mlir::dyn_cast<fir::BaseBoxType>(load.getType())) {3364 // fir.box is a special case because it is considered an ssa value in3365 // fir, but it is lowered as a pointer to a descriptor. So3366 // fir.ref<fir.box> and fir.box end up being the same llvm types and3367 // loading a fir.ref<fir.box> is implemented as taking a snapshot of the3368 // descriptor value into a new descriptor temp.3369 auto inputBoxStorage = adaptor.getOperands()[0];3370 mlir::Value newBoxStorage;3371 mlir::Location loc = load.getLoc();3372 if (auto callOp = mlir::dyn_cast_or_null<mlir::LLVM::CallOp>(3373 inputBoxStorage.getDefiningOp())) {3374 if (callOp.getCallee() &&3375 ((*callOp.getCallee())3376 .starts_with(RTNAME_STRING(CUFAllocDescriptor)) ||3377 (*callOp.getCallee()).starts_with("__tgt_acc_get_deviceptr"))) {3378 // CUDA Fortran local descriptor are allocated in managed memory. So3379 // new storage must be allocated the same way.3380 auto mod = load->getParentOfType<mlir::ModuleOp>();3381 newBoxStorage =3382 genCUFAllocDescriptor(loc, rewriter, mod, boxTy, lowerTy());3383 }3384 }3385 if (!newBoxStorage)3386 newBoxStorage = genAllocaAndAddrCastWithType(loc, llvmLoadTy,3387 defaultAlign, rewriter);3388 3389 TypePair boxTypePair{boxTy, llvmLoadTy};3390 mlir::Value boxSize =3391 computeBoxSize(loc, boxTypePair, inputBoxStorage, rewriter);3392 auto memcpy = mlir::LLVM::MemcpyOp::create(3393 rewriter, loc, newBoxStorage, inputBoxStorage, boxSize, isVolatile);3394 3395 if (std::optional<mlir::ArrayAttr> optionalTag = load.getTbaa())3396 memcpy.setTBAATags(*optionalTag);3397 else3398 attachTBAATag(memcpy, boxTy, boxTy, nullptr);3399 rewriter.replaceOp(load, newBoxStorage);3400 } else {3401 mlir::LLVM::LoadOp loadOp =3402 mlir::LLVM::LoadOp::create(rewriter, load.getLoc(), llvmLoadTy,3403 adaptor.getOperands(), load->getAttrs());3404 loadOp.setVolatile_(isVolatile);3405 if (std::optional<mlir::ArrayAttr> optionalTag = load.getTbaa())3406 loadOp.setTBAATags(*optionalTag);3407 else3408 attachTBAATag(loadOp, load.getType(), load.getType(), nullptr);3409 if (std::optional<mlir::ArrayAttr> optionalAccessGroups =3410 load.getAccessGroups())3411 loadOp.setAccessGroups(*optionalAccessGroups);3412 rewriter.replaceOp(load, loadOp.getResult());3413 }3414 return mlir::success();3415 }3416};3417 3418template <typename OpTy>3419struct DoConcurrentSpecifierOpConversion : public fir::FIROpConversion<OpTy> {3420 using fir::FIROpConversion<OpTy>::FIROpConversion;3421 llvm::LogicalResult3422 matchAndRewrite(OpTy specifier, typename OpTy::Adaptor adaptor,3423 mlir::ConversionPatternRewriter &rewriter) const override {3424#ifdef EXPENSIVE_CHECKS3425 auto uses = mlir::SymbolTable::getSymbolUses(3426 specifier, specifier->template getParentOfType<mlir::ModuleOp>());3427 3428 // `fir.local|fir.declare_reduction` ops are not supposed to have any uses3429 // at this point (i.e. during lowering to LLVM). In case of serialization,3430 // the `fir.do_concurrent` users are expected to have been lowered to3431 // `fir.do_loop` nests. In case of parallelization, the `fir.do_concurrent`3432 // users are expected to have been lowered to the target parallel model3433 // (e.g. OpenMP).3434 assert(uses && uses->empty());3435#endif3436 3437 rewriter.eraseOp(specifier);3438 return mlir::success();3439 }3440};3441 3442/// Lower `fir.no_reassoc` to LLVM IR dialect.3443/// TODO: how do we want to enforce this in LLVM-IR? Can we manipulate the fast3444/// math flags?3445struct NoReassocOpConversion : public fir::FIROpConversion<fir::NoReassocOp> {3446 using FIROpConversion::FIROpConversion;3447 3448 llvm::LogicalResult3449 matchAndRewrite(fir::NoReassocOp noreassoc, OpAdaptor adaptor,3450 mlir::ConversionPatternRewriter &rewriter) const override {3451 rewriter.replaceOp(noreassoc, adaptor.getOperands()[0]);3452 return mlir::success();3453 }3454};3455 3456static void genCondBrOp(mlir::Location loc, mlir::Value cmp, mlir::Block *dest,3457 std::optional<mlir::ValueRange> destOps,3458 mlir::ConversionPatternRewriter &rewriter,3459 mlir::Block *newBlock) {3460 if (destOps)3461 mlir::LLVM::CondBrOp::create(rewriter, loc, cmp, dest, *destOps, newBlock,3462 mlir::ValueRange());3463 else3464 mlir::LLVM::CondBrOp::create(rewriter, loc, cmp, dest, newBlock);3465}3466 3467template <typename A, typename B>3468static void genBrOp(A caseOp, mlir::Block *dest, std::optional<B> destOps,3469 mlir::ConversionPatternRewriter &rewriter) {3470 if (destOps)3471 rewriter.replaceOpWithNewOp<mlir::LLVM::BrOp>(caseOp, *destOps, dest);3472 else3473 rewriter.replaceOpWithNewOp<mlir::LLVM::BrOp>(caseOp, B{}, dest);3474}3475 3476static void genCaseLadderStep(mlir::Location loc, mlir::Value cmp,3477 mlir::Block *dest,3478 std::optional<mlir::ValueRange> destOps,3479 mlir::ConversionPatternRewriter &rewriter) {3480 auto *thisBlock = rewriter.getInsertionBlock();3481 auto *newBlock = createBlock(rewriter, dest);3482 rewriter.setInsertionPointToEnd(thisBlock);3483 genCondBrOp(loc, cmp, dest, destOps, rewriter, newBlock);3484 rewriter.setInsertionPointToEnd(newBlock);3485}3486 3487/// Conversion of `fir.select_case`3488///3489/// The `fir.select_case` operation is converted to a if-then-else ladder.3490/// Depending on the case condition type, one or several comparison and3491/// conditional branching can be generated.3492///3493/// A point value case such as `case(4)`, a lower bound case such as3494/// `case(5:)` or an upper bound case such as `case(:3)` are converted to a3495/// simple comparison between the selector value and the constant value in the3496/// case. The block associated with the case condition is then executed if3497/// the comparison succeed otherwise it branch to the next block with the3498/// comparison for the next case conditon.3499///3500/// A closed interval case condition such as `case(7:10)` is converted with a3501/// first comparison and conditional branching for the lower bound. If3502/// successful, it branch to a second block with the comparison for the3503/// upper bound in the same case condition.3504///3505/// TODO: lowering of CHARACTER type cases is not handled yet.3506struct SelectCaseOpConversion : public fir::FIROpConversion<fir::SelectCaseOp> {3507 using FIROpConversion::FIROpConversion;3508 3509 llvm::LogicalResult3510 matchAndRewrite(fir::SelectCaseOp caseOp, OpAdaptor adaptor,3511 mlir::ConversionPatternRewriter &rewriter) const override {3512 unsigned conds = caseOp.getNumConditions();3513 llvm::ArrayRef<mlir::Attribute> cases = caseOp.getCases().getValue();3514 // Type can be CHARACTER, INTEGER, or LOGICAL (C1145)3515 auto ty = caseOp.getSelector().getType();3516 if (mlir::isa<fir::CharacterType>(ty)) {3517 TODO(caseOp.getLoc(), "fir.select_case codegen with character type");3518 return mlir::failure();3519 }3520 mlir::Value selector = caseOp.getSelector(adaptor.getOperands());3521 auto loc = caseOp.getLoc();3522 for (unsigned t = 0; t != conds; ++t) {3523 mlir::Block *dest = caseOp.getSuccessor(t);3524 std::optional<mlir::ValueRange> destOps =3525 caseOp.getSuccessorOperands(adaptor.getOperands(), t);3526 std::optional<mlir::ValueRange> cmpOps =3527 *caseOp.getCompareOperands(adaptor.getOperands(), t);3528 mlir::Attribute attr = cases[t];3529 assert(mlir::isa<mlir::UnitAttr>(attr) || cmpOps.has_value());3530 if (mlir::isa<fir::PointIntervalAttr>(attr)) {3531 auto cmp = mlir::LLVM::ICmpOp::create(rewriter, loc,3532 mlir::LLVM::ICmpPredicate::eq,3533 selector, cmpOps->front());3534 genCaseLadderStep(loc, cmp, dest, destOps, rewriter);3535 continue;3536 }3537 if (mlir::isa<fir::LowerBoundAttr>(attr)) {3538 auto cmp = mlir::LLVM::ICmpOp::create(rewriter, loc,3539 mlir::LLVM::ICmpPredicate::sle,3540 cmpOps->front(), selector);3541 genCaseLadderStep(loc, cmp, dest, destOps, rewriter);3542 continue;3543 }3544 if (mlir::isa<fir::UpperBoundAttr>(attr)) {3545 auto cmp = mlir::LLVM::ICmpOp::create(rewriter, loc,3546 mlir::LLVM::ICmpPredicate::sle,3547 selector, cmpOps->front());3548 genCaseLadderStep(loc, cmp, dest, destOps, rewriter);3549 continue;3550 }3551 if (mlir::isa<fir::ClosedIntervalAttr>(attr)) {3552 mlir::Value caseArg0 = *cmpOps->begin();3553 auto cmp0 = mlir::LLVM::ICmpOp::create(3554 rewriter, loc, mlir::LLVM::ICmpPredicate::sle, caseArg0, selector);3555 auto *thisBlock = rewriter.getInsertionBlock();3556 auto *newBlock1 = createBlock(rewriter, dest);3557 auto *newBlock2 = createBlock(rewriter, dest);3558 rewriter.setInsertionPointToEnd(thisBlock);3559 mlir::LLVM::CondBrOp::create(rewriter, loc, cmp0, newBlock1, newBlock2);3560 rewriter.setInsertionPointToEnd(newBlock1);3561 mlir::Value caseArg1 = *(cmpOps->begin() + 1);3562 auto cmp1 = mlir::LLVM::ICmpOp::create(3563 rewriter, loc, mlir::LLVM::ICmpPredicate::sle, selector, caseArg1);3564 genCondBrOp(loc, cmp1, dest, destOps, rewriter, newBlock2);3565 rewriter.setInsertionPointToEnd(newBlock2);3566 continue;3567 }3568 assert(mlir::isa<mlir::UnitAttr>(attr));3569 assert((t + 1 == conds) && "unit must be last");3570 genBrOp(caseOp, dest, destOps, rewriter);3571 }3572 return mlir::success();3573 }3574};3575 3576/// Base class for SelectOpConversion and SelectRankOpConversion.3577template <typename OP>3578struct SelectOpConversionBase : public fir::FIROpConversion<OP> {3579 using fir::FIROpConversion<OP>::FIROpConversion;3580 3581private:3582 /// Helper function for converting select ops. This function converts the3583 /// signature of the given block. If the new block signature is different from3584 /// `expectedTypes`, returns "failure".3585 llvm::FailureOr<mlir::Block *>3586 getConvertedBlock(mlir::ConversionPatternRewriter &rewriter,3587 mlir::Operation *branchOp, mlir::Block *block,3588 mlir::TypeRange expectedTypes) const {3589 const mlir::TypeConverter *converter = this->getTypeConverter();3590 assert(converter && "expected non-null type converter");3591 assert(!block->isEntryBlock() && "entry blocks have no predecessors");3592 3593 // There is nothing to do if the types already match.3594 if (block->getArgumentTypes() == expectedTypes)3595 return block;3596 3597 // Compute the new block argument types and convert the block.3598 std::optional<mlir::TypeConverter::SignatureConversion> conversion =3599 converter->convertBlockSignature(block);3600 if (!conversion)3601 return rewriter.notifyMatchFailure(branchOp,3602 "could not compute block signature");3603 if (expectedTypes != conversion->getConvertedTypes())3604 return rewriter.notifyMatchFailure(branchOp,3605 "mismatch between adaptor operand "3606 "types and computed block signature");3607 return rewriter.applySignatureConversion(block, *conversion, converter);3608 }3609 3610protected:3611 llvm::LogicalResult3612 selectMatchAndRewrite(OP select, typename OP::Adaptor adaptor,3613 mlir::ConversionPatternRewriter &rewriter) const {3614 unsigned conds = select.getNumConditions();3615 auto cases = select.getCases().getValue();3616 mlir::Value selector = adaptor.getSelector();3617 auto loc = select.getLoc();3618 assert(conds > 0 && "select must have cases");3619 3620 llvm::SmallVector<mlir::Block *> destinations;3621 llvm::SmallVector<mlir::ValueRange> destinationsOperands;3622 mlir::Block *defaultDestination;3623 mlir::ValueRange defaultOperands;3624 // LLVM::SwitchOp selector type and the case values types3625 // must have the same bit width, so cast the selector to i64,3626 // and use i64 for the case values. It is hard to imagine3627 // a computed GO TO with the number of labels in the label-list3628 // bigger than INT_MAX, but let's use i64 to be on the safe side.3629 // Moreover, fir.select operation is more relaxed than3630 // a Fortran computed GO TO, so it may specify such a case value3631 // even if there is just a single label/case.3632 llvm::SmallVector<int64_t> caseValues;3633 3634 for (unsigned t = 0; t != conds; ++t) {3635 mlir::Block *dest = select.getSuccessor(t);3636 auto destOps = select.getSuccessorOperands(adaptor.getOperands(), t);3637 const mlir::Attribute &attr = cases[t];3638 if (auto intAttr = mlir::dyn_cast<mlir::IntegerAttr>(attr)) {3639 destinationsOperands.push_back(destOps ? *destOps : mlir::ValueRange{});3640 auto convertedBlock =3641 getConvertedBlock(rewriter, select, dest,3642 mlir::TypeRange(destinationsOperands.back()));3643 if (mlir::failed(convertedBlock))3644 return mlir::failure();3645 destinations.push_back(*convertedBlock);3646 caseValues.push_back(intAttr.getInt());3647 continue;3648 }3649 assert(mlir::dyn_cast_or_null<mlir::UnitAttr>(attr));3650 assert((t + 1 == conds) && "unit must be last");3651 defaultOperands = destOps ? *destOps : mlir::ValueRange{};3652 auto convertedBlock = getConvertedBlock(rewriter, select, dest,3653 mlir::TypeRange(defaultOperands));3654 if (mlir::failed(convertedBlock))3655 return mlir::failure();3656 defaultDestination = *convertedBlock;3657 }3658 3659 selector =3660 this->integerCast(loc, rewriter, rewriter.getI64Type(), selector);3661 3662 rewriter.replaceOpWithNewOp<mlir::LLVM::SwitchOp>(3663 select, selector,3664 /*defaultDestination=*/defaultDestination,3665 /*defaultOperands=*/defaultOperands,3666 /*caseValues=*/rewriter.getI64VectorAttr(caseValues),3667 /*caseDestinations=*/destinations,3668 /*caseOperands=*/destinationsOperands,3669 /*branchWeights=*/llvm::ArrayRef<std::int32_t>());3670 return mlir::success();3671 }3672};3673/// conversion of fir::SelectOp to an if-then-else ladder3674struct SelectOpConversion : public SelectOpConversionBase<fir::SelectOp> {3675 using SelectOpConversionBase::SelectOpConversionBase;3676 3677 llvm::LogicalResult3678 matchAndRewrite(fir::SelectOp op, OpAdaptor adaptor,3679 mlir::ConversionPatternRewriter &rewriter) const override {3680 return this->selectMatchAndRewrite(op, adaptor, rewriter);3681 }3682};3683 3684/// conversion of fir::SelectRankOp to an if-then-else ladder3685struct SelectRankOpConversion3686 : public SelectOpConversionBase<fir::SelectRankOp> {3687 using SelectOpConversionBase::SelectOpConversionBase;3688 3689 llvm::LogicalResult3690 matchAndRewrite(fir::SelectRankOp op, OpAdaptor adaptor,3691 mlir::ConversionPatternRewriter &rewriter) const override {3692 return this->selectMatchAndRewrite(op, adaptor, rewriter);3693 }3694};3695 3696/// Lower `fir.select_type` to LLVM IR dialect.3697struct SelectTypeOpConversion : public fir::FIROpConversion<fir::SelectTypeOp> {3698 using FIROpConversion::FIROpConversion;3699 3700 llvm::LogicalResult3701 matchAndRewrite(fir::SelectTypeOp select, OpAdaptor adaptor,3702 mlir::ConversionPatternRewriter &rewriter) const override {3703 mlir::emitError(select.getLoc(),3704 "fir.select_type should have already been converted");3705 return mlir::failure();3706 }3707};3708 3709/// `fir.store` --> `llvm.store`3710struct StoreOpConversion : public fir::FIROpConversion<fir::StoreOp> {3711 using FIROpConversion::FIROpConversion;3712 3713 llvm::LogicalResult3714 matchAndRewrite(fir::StoreOp store, OpAdaptor adaptor,3715 mlir::ConversionPatternRewriter &rewriter) const override {3716 mlir::Location loc = store.getLoc();3717 mlir::Type storeTy = store.getValue().getType();3718 mlir::Value llvmValue = adaptor.getValue();3719 mlir::Value llvmMemref = adaptor.getMemref();3720 mlir::LLVM::AliasAnalysisOpInterface newOp;3721 const bool isVolatile =3722 fir::isa_volatile_type(store.getMemref().getType()) ||3723 fir::isa_volatile_type(store.getValue().getType());3724 if (auto boxTy = mlir::dyn_cast<fir::BaseBoxType>(storeTy)) {3725 mlir::Type llvmBoxTy = lowerTy().convertBoxTypeAsStruct(boxTy);3726 // Always use memcpy because LLVM is not as effective at optimizing3727 // aggregate loads/stores as it is optimizing memcpy.3728 TypePair boxTypePair{boxTy, llvmBoxTy};3729 mlir::Value boxSize =3730 computeBoxSize(loc, boxTypePair, llvmValue, rewriter);3731 newOp = mlir::LLVM::MemcpyOp::create(rewriter, loc, llvmMemref, llvmValue,3732 boxSize, isVolatile);3733 } else {3734 mlir::LLVM::StoreOp storeOp =3735 mlir::LLVM::StoreOp::create(rewriter, loc, llvmValue, llvmMemref);3736 3737 if (isVolatile)3738 storeOp.setVolatile_(true);3739 3740 if (store.getNontemporal())3741 storeOp.setNontemporal(true);3742 3743 if (std::optional<mlir::ArrayAttr> optionalAccessGroups =3744 store.getAccessGroups())3745 storeOp.setAccessGroups(*optionalAccessGroups);3746 3747 newOp = storeOp;3748 }3749 if (std::optional<mlir::ArrayAttr> optionalTag = store.getTbaa())3750 newOp.setTBAATags(*optionalTag);3751 else3752 attachTBAATag(newOp, storeTy, storeTy, nullptr);3753 rewriter.eraseOp(store);3754 return mlir::success();3755 }3756};3757 3758/// `fir.copy` --> `llvm.memcpy` or `llvm.memmove`3759struct CopyOpConversion : public fir::FIROpConversion<fir::CopyOp> {3760 using FIROpConversion::FIROpConversion;3761 3762 llvm::LogicalResult3763 matchAndRewrite(fir::CopyOp copy, OpAdaptor adaptor,3764 mlir::ConversionPatternRewriter &rewriter) const override {3765 mlir::Location loc = copy.getLoc();3766 const bool isVolatile =3767 fir::isa_volatile_type(copy.getSource().getType()) ||3768 fir::isa_volatile_type(copy.getDestination().getType());3769 mlir::Value llvmSource = adaptor.getSource();3770 mlir::Value llvmDestination = adaptor.getDestination();3771 mlir::Type i64Ty = mlir::IntegerType::get(rewriter.getContext(), 64);3772 mlir::Type copyTy = fir::unwrapRefType(copy.getSource().getType());3773 mlir::Value copySize = genTypeStrideInBytes(3774 loc, i64Ty, rewriter, convertType(copyTy), getDataLayout());3775 3776 mlir::LLVM::AliasAnalysisOpInterface newOp;3777 if (copy.getNoOverlap())3778 newOp = mlir::LLVM::MemcpyOp::create(rewriter, loc, llvmDestination,3779 llvmSource, copySize, isVolatile);3780 else3781 newOp = mlir::LLVM::MemmoveOp::create(rewriter, loc, llvmDestination,3782 llvmSource, copySize, isVolatile);3783 3784 // TODO: propagate TBAA once FirAliasTagOpInterface added to CopyOp.3785 attachTBAATag(newOp, copyTy, copyTy, nullptr);3786 rewriter.eraseOp(copy);3787 return mlir::success();3788 }3789};3790 3791namespace {3792 3793/// Convert `fir.unboxchar` into two `llvm.extractvalue` instructions. One for3794/// the character buffer and one for the buffer length.3795struct UnboxCharOpConversion : public fir::FIROpConversion<fir::UnboxCharOp> {3796 using FIROpConversion::FIROpConversion;3797 3798 llvm::LogicalResult3799 matchAndRewrite(fir::UnboxCharOp unboxchar, OpAdaptor adaptor,3800 mlir::ConversionPatternRewriter &rewriter) const override {3801 mlir::Type lenTy = convertType(unboxchar.getType(1));3802 mlir::Value tuple = adaptor.getOperands()[0];3803 3804 mlir::Location loc = unboxchar.getLoc();3805 mlir::Value ptrToBuffer =3806 mlir::LLVM::ExtractValueOp::create(rewriter, loc, tuple, 0);3807 3808 auto len = mlir::LLVM::ExtractValueOp::create(rewriter, loc, tuple, 1);3809 mlir::Value lenAfterCast = integerCast(loc, rewriter, lenTy, len);3810 3811 rewriter.replaceOp(unboxchar,3812 llvm::ArrayRef<mlir::Value>{ptrToBuffer, lenAfterCast});3813 return mlir::success();3814 }3815};3816 3817/// Lower `fir.unboxproc` operation. Unbox a procedure box value, yielding its3818/// components.3819/// TODO: Part of supporting Fortran 2003 procedure pointers.3820struct UnboxProcOpConversion : public fir::FIROpConversion<fir::UnboxProcOp> {3821 using FIROpConversion::FIROpConversion;3822 3823 llvm::LogicalResult3824 matchAndRewrite(fir::UnboxProcOp unboxproc, OpAdaptor adaptor,3825 mlir::ConversionPatternRewriter &rewriter) const override {3826 TODO(unboxproc.getLoc(), "fir.unboxproc codegen");3827 return mlir::failure();3828 }3829};3830 3831/// convert to LLVM IR dialect `undef`3832struct UndefOpConversion : public fir::FIROpConversion<fir::UndefOp> {3833 using FIROpConversion::FIROpConversion;3834 3835 llvm::LogicalResult3836 matchAndRewrite(fir::UndefOp undef, OpAdaptor,3837 mlir::ConversionPatternRewriter &rewriter) const override {3838 if (mlir::isa<fir::DummyScopeType>(undef.getType())) {3839 // Dummy scoping is used for Fortran analyses like AA. Once it gets to3840 // pre-codegen rewrite it is erased and a fir.undef is created to3841 // feed to the fir declare operation. Thus, during codegen, we can3842 // simply erase is as it is no longer used.3843 rewriter.eraseOp(undef);3844 return mlir::success();3845 }3846 rewriter.replaceOpWithNewOp<mlir::LLVM::UndefOp>(3847 undef, convertType(undef.getType()));3848 return mlir::success();3849 }3850};3851 3852struct ZeroOpConversion : public fir::FIROpConversion<fir::ZeroOp> {3853 using FIROpConversion::FIROpConversion;3854 3855 llvm::LogicalResult3856 matchAndRewrite(fir::ZeroOp zero, OpAdaptor,3857 mlir::ConversionPatternRewriter &rewriter) const override {3858 mlir::Type ty = convertType(zero.getType());3859 rewriter.replaceOpWithNewOp<mlir::LLVM::ZeroOp>(zero, ty);3860 return mlir::success();3861 }3862};3863 3864/// `fir.unreachable` --> `llvm.unreachable`3865struct UnreachableOpConversion3866 : public fir::FIROpConversion<fir::UnreachableOp> {3867 using FIROpConversion::FIROpConversion;3868 3869 llvm::LogicalResult3870 matchAndRewrite(fir::UnreachableOp unreach, OpAdaptor adaptor,3871 mlir::ConversionPatternRewriter &rewriter) const override {3872 rewriter.replaceOpWithNewOp<mlir::LLVM::UnreachableOp>(unreach);3873 return mlir::success();3874 }3875};3876 3877/// `fir.is_present` -->3878/// ```3879/// %0 = llvm.mlir.constant(0 : i64)3880/// %1 = llvm.ptrtoint %03881/// %2 = llvm.icmp "ne" %1, %0 : i643882/// ```3883struct IsPresentOpConversion : public fir::FIROpConversion<fir::IsPresentOp> {3884 using FIROpConversion::FIROpConversion;3885 3886 llvm::LogicalResult3887 matchAndRewrite(fir::IsPresentOp isPresent, OpAdaptor adaptor,3888 mlir::ConversionPatternRewriter &rewriter) const override {3889 mlir::Type idxTy = lowerTy().indexType();3890 mlir::Location loc = isPresent.getLoc();3891 auto ptr = adaptor.getOperands()[0];3892 3893 if (mlir::isa<fir::BoxCharType>(isPresent.getVal().getType())) {3894 [[maybe_unused]] auto structTy =3895 mlir::cast<mlir::LLVM::LLVMStructType>(ptr.getType());3896 assert(!structTy.isOpaque() && !structTy.getBody().empty());3897 3898 ptr = mlir::LLVM::ExtractValueOp::create(rewriter, loc, ptr, 0);3899 }3900 mlir::LLVM::ConstantOp c0 =3901 fir::genConstantIndex(isPresent.getLoc(), idxTy, rewriter, 0);3902 auto addr = mlir::LLVM::PtrToIntOp::create(rewriter, loc, idxTy, ptr);3903 rewriter.replaceOpWithNewOp<mlir::LLVM::ICmpOp>(3904 isPresent, mlir::LLVM::ICmpPredicate::ne, addr, c0);3905 3906 return mlir::success();3907 }3908};3909 3910/// Create value signaling an absent optional argument in a call, e.g.3911/// `fir.absent !fir.ref<i64>` --> `llvm.mlir.zero : !llvm.ptr<i64>`3912struct AbsentOpConversion : public fir::FIROpConversion<fir::AbsentOp> {3913 using FIROpConversion::FIROpConversion;3914 3915 llvm::LogicalResult3916 matchAndRewrite(fir::AbsentOp absent, OpAdaptor,3917 mlir::ConversionPatternRewriter &rewriter) const override {3918 mlir::Type ty = convertType(absent.getType());3919 rewriter.replaceOpWithNewOp<mlir::LLVM::ZeroOp>(absent, ty);3920 return mlir::success();3921 }3922};3923 3924//3925// Primitive operations on Complex types3926//3927 3928template <typename OPTY>3929static inline mlir::LLVM::FastmathFlagsAttr getLLVMFMFAttr(OPTY op) {3930 return mlir::LLVM::FastmathFlagsAttr::get(3931 op.getContext(),3932 mlir::arith::convertArithFastMathFlagsToLLVM(op.getFastmath()));3933}3934 3935/// Generate inline code for complex addition/subtraction3936template <typename LLVMOP, typename OPTY>3937static mlir::LLVM::InsertValueOp3938complexSum(OPTY sumop, mlir::ValueRange opnds,3939 mlir::ConversionPatternRewriter &rewriter,3940 const fir::LLVMTypeConverter &lowering) {3941 mlir::LLVM::FastmathFlagsAttr fmf = getLLVMFMFAttr(sumop);3942 mlir::Value a = opnds[0];3943 mlir::Value b = opnds[1];3944 auto loc = sumop.getLoc();3945 mlir::Type eleTy = lowering.convertType(getComplexEleTy(sumop.getType()));3946 mlir::Type ty = lowering.convertType(sumop.getType());3947 auto x0 = mlir::LLVM::ExtractValueOp::create(rewriter, loc, a, 0);3948 auto y0 = mlir::LLVM::ExtractValueOp::create(rewriter, loc, a, 1);3949 auto x1 = mlir::LLVM::ExtractValueOp::create(rewriter, loc, b, 0);3950 auto y1 = mlir::LLVM::ExtractValueOp::create(rewriter, loc, b, 1);3951 auto rx = LLVMOP::create(rewriter, loc, eleTy, x0, x1, fmf);3952 auto ry = LLVMOP::create(rewriter, loc, eleTy, y0, y1, fmf);3953 auto r0 = mlir::LLVM::UndefOp::create(rewriter, loc, ty);3954 llvm::SmallVector<int64_t> pos{0};3955 auto r1 = mlir::LLVM::InsertValueOp::create(rewriter, loc, r0, rx, pos);3956 return mlir::LLVM::InsertValueOp::create(rewriter, loc, r1, ry, 1);3957}3958} // namespace3959 3960namespace {3961struct AddcOpConversion : public fir::FIROpConversion<fir::AddcOp> {3962 using FIROpConversion::FIROpConversion;3963 3964 llvm::LogicalResult3965 matchAndRewrite(fir::AddcOp addc, OpAdaptor adaptor,3966 mlir::ConversionPatternRewriter &rewriter) const override {3967 // given: (x + iy) + (x' + iy')3968 // result: (x + x') + i(y + y')3969 auto r = complexSum<mlir::LLVM::FAddOp>(addc, adaptor.getOperands(),3970 rewriter, lowerTy());3971 rewriter.replaceOp(addc, r.getResult());3972 return mlir::success();3973 }3974};3975 3976struct SubcOpConversion : public fir::FIROpConversion<fir::SubcOp> {3977 using FIROpConversion::FIROpConversion;3978 3979 llvm::LogicalResult3980 matchAndRewrite(fir::SubcOp subc, OpAdaptor adaptor,3981 mlir::ConversionPatternRewriter &rewriter) const override {3982 // given: (x + iy) - (x' + iy')3983 // result: (x - x') + i(y - y')3984 auto r = complexSum<mlir::LLVM::FSubOp>(subc, adaptor.getOperands(),3985 rewriter, lowerTy());3986 rewriter.replaceOp(subc, r.getResult());3987 return mlir::success();3988 }3989};3990 3991/// Inlined complex multiply3992struct MulcOpConversion : public fir::FIROpConversion<fir::MulcOp> {3993 using FIROpConversion::FIROpConversion;3994 3995 llvm::LogicalResult3996 matchAndRewrite(fir::MulcOp mulc, OpAdaptor adaptor,3997 mlir::ConversionPatternRewriter &rewriter) const override {3998 // TODO: Can we use a call to __muldc3 ?3999 // given: (x + iy) * (x' + iy')4000 // result: (xx'-yy')+i(xy'+yx')4001 mlir::LLVM::FastmathFlagsAttr fmf = getLLVMFMFAttr(mulc);4002 mlir::Value a = adaptor.getOperands()[0];4003 mlir::Value b = adaptor.getOperands()[1];4004 auto loc = mulc.getLoc();4005 mlir::Type eleTy = convertType(getComplexEleTy(mulc.getType()));4006 mlir::Type ty = convertType(mulc.getType());4007 auto x0 = mlir::LLVM::ExtractValueOp::create(rewriter, loc, a, 0);4008 auto y0 = mlir::LLVM::ExtractValueOp::create(rewriter, loc, a, 1);4009 auto x1 = mlir::LLVM::ExtractValueOp::create(rewriter, loc, b, 0);4010 auto y1 = mlir::LLVM::ExtractValueOp::create(rewriter, loc, b, 1);4011 auto xx = mlir::LLVM::FMulOp::create(rewriter, loc, eleTy, x0, x1, fmf);4012 auto yx = mlir::LLVM::FMulOp::create(rewriter, loc, eleTy, y0, x1, fmf);4013 auto xy = mlir::LLVM::FMulOp::create(rewriter, loc, eleTy, x0, y1, fmf);4014 auto ri = mlir::LLVM::FAddOp::create(rewriter, loc, eleTy, xy, yx, fmf);4015 auto yy = mlir::LLVM::FMulOp::create(rewriter, loc, eleTy, y0, y1, fmf);4016 auto rr = mlir::LLVM::FSubOp::create(rewriter, loc, eleTy, xx, yy, fmf);4017 auto ra = mlir::LLVM::UndefOp::create(rewriter, loc, ty);4018 llvm::SmallVector<int64_t> pos{0};4019 auto r1 = mlir::LLVM::InsertValueOp::create(rewriter, loc, ra, rr, pos);4020 auto r0 = mlir::LLVM::InsertValueOp::create(rewriter, loc, r1, ri, 1);4021 rewriter.replaceOp(mulc, r0.getResult());4022 return mlir::success();4023 }4024};4025 4026/// Inlined complex division4027struct DivcOpConversion : public fir::FIROpConversion<fir::DivcOp> {4028 using FIROpConversion::FIROpConversion;4029 4030 llvm::LogicalResult4031 matchAndRewrite(fir::DivcOp divc, OpAdaptor adaptor,4032 mlir::ConversionPatternRewriter &rewriter) const override {4033 // TODO: Can we use a call to __divdc3 instead?4034 // Just generate inline code for now.4035 // given: (x + iy) / (x' + iy')4036 // result: ((xx'+yy')/d) + i((yx'-xy')/d) where d = x'x' + y'y'4037 mlir::LLVM::FastmathFlagsAttr fmf = getLLVMFMFAttr(divc);4038 mlir::Value a = adaptor.getOperands()[0];4039 mlir::Value b = adaptor.getOperands()[1];4040 auto loc = divc.getLoc();4041 mlir::Type eleTy = convertType(getComplexEleTy(divc.getType()));4042 mlir::Type ty = convertType(divc.getType());4043 auto x0 = mlir::LLVM::ExtractValueOp::create(rewriter, loc, a, 0);4044 auto y0 = mlir::LLVM::ExtractValueOp::create(rewriter, loc, a, 1);4045 auto x1 = mlir::LLVM::ExtractValueOp::create(rewriter, loc, b, 0);4046 auto y1 = mlir::LLVM::ExtractValueOp::create(rewriter, loc, b, 1);4047 auto xx = mlir::LLVM::FMulOp::create(rewriter, loc, eleTy, x0, x1, fmf);4048 auto x1x1 = mlir::LLVM::FMulOp::create(rewriter, loc, eleTy, x1, x1, fmf);4049 auto yx = mlir::LLVM::FMulOp::create(rewriter, loc, eleTy, y0, x1, fmf);4050 auto xy = mlir::LLVM::FMulOp::create(rewriter, loc, eleTy, x0, y1, fmf);4051 auto yy = mlir::LLVM::FMulOp::create(rewriter, loc, eleTy, y0, y1, fmf);4052 auto y1y1 = mlir::LLVM::FMulOp::create(rewriter, loc, eleTy, y1, y1, fmf);4053 auto d = mlir::LLVM::FAddOp::create(rewriter, loc, eleTy, x1x1, y1y1, fmf);4054 auto rrn = mlir::LLVM::FAddOp::create(rewriter, loc, eleTy, xx, yy, fmf);4055 auto rin = mlir::LLVM::FSubOp::create(rewriter, loc, eleTy, yx, xy, fmf);4056 auto rr = mlir::LLVM::FDivOp::create(rewriter, loc, eleTy, rrn, d, fmf);4057 auto ri = mlir::LLVM::FDivOp::create(rewriter, loc, eleTy, rin, d, fmf);4058 auto ra = mlir::LLVM::UndefOp::create(rewriter, loc, ty);4059 llvm::SmallVector<int64_t> pos{0};4060 auto r1 = mlir::LLVM::InsertValueOp::create(rewriter, loc, ra, rr, pos);4061 auto r0 = mlir::LLVM::InsertValueOp::create(rewriter, loc, r1, ri, 1);4062 rewriter.replaceOp(divc, r0.getResult());4063 return mlir::success();4064 }4065};4066 4067/// Inlined complex negation4068struct NegcOpConversion : public fir::FIROpConversion<fir::NegcOp> {4069 using FIROpConversion::FIROpConversion;4070 4071 llvm::LogicalResult4072 matchAndRewrite(fir::NegcOp neg, OpAdaptor adaptor,4073 mlir::ConversionPatternRewriter &rewriter) const override {4074 // given: -(x + iy)4075 // result: -x - iy4076 auto eleTy = convertType(getComplexEleTy(neg.getType()));4077 auto loc = neg.getLoc();4078 mlir::Value o0 = adaptor.getOperands()[0];4079 auto rp = mlir::LLVM::ExtractValueOp::create(rewriter, loc, o0, 0);4080 auto ip = mlir::LLVM::ExtractValueOp::create(rewriter, loc, o0, 1);4081 auto nrp = mlir::LLVM::FNegOp::create(rewriter, loc, eleTy, rp);4082 auto nip = mlir::LLVM::FNegOp::create(rewriter, loc, eleTy, ip);4083 llvm::SmallVector<int64_t> pos{0};4084 auto r = mlir::LLVM::InsertValueOp::create(rewriter, loc, o0, nrp, pos);4085 rewriter.replaceOpWithNewOp<mlir::LLVM::InsertValueOp>(neg, r, nip, 1);4086 return mlir::success();4087 }4088};4089 4090struct BoxOffsetOpConversion : public fir::FIROpConversion<fir::BoxOffsetOp> {4091 using FIROpConversion::FIROpConversion;4092 4093 llvm::LogicalResult4094 matchAndRewrite(fir::BoxOffsetOp boxOffset, OpAdaptor adaptor,4095 mlir::ConversionPatternRewriter &rewriter) const override {4096 4097 mlir::Type pty = ::getLlvmPtrType(boxOffset.getContext());4098 mlir::Type boxRefType = fir::unwrapRefType(boxOffset.getBoxRef().getType());4099 4100 assert((mlir::isa<fir::BaseBoxType>(boxRefType) ||4101 mlir::isa<fir::BoxCharType>(boxRefType)) &&4102 "boxRef should be a reference to either fir.box or fir.boxchar");4103 4104 mlir::Type llvmBoxTy;4105 int fieldId;4106 if (auto boxType = mlir::dyn_cast_or_null<fir::BaseBoxType>(boxRefType)) {4107 llvmBoxTy = lowerTy().convertBoxTypeAsStruct(4108 mlir::cast<fir::BaseBoxType>(boxType));4109 fieldId = boxOffset.getField() == fir::BoxFieldAttr::derived_type4110 ? getTypeDescFieldId(boxType)4111 : kAddrPosInBox;4112 } else {4113 auto boxCharType = mlir::cast<fir::BoxCharType>(boxRefType);4114 llvmBoxTy = lowerTy().convertType(boxCharType);4115 fieldId = kAddrPosInBox;4116 }4117 rewriter.replaceOpWithNewOp<mlir::LLVM::GEPOp>(4118 boxOffset, pty, llvmBoxTy, adaptor.getBoxRef(),4119 llvm::ArrayRef<mlir::LLVM::GEPArg>{0, fieldId});4120 return mlir::success();4121 }4122};4123 4124/// Conversion pattern for operation that must be dead. The information in these4125/// operations is used by other operation. At this point they should not have4126/// anymore uses.4127/// These operations are normally dead after the pre-codegen pass.4128template <typename FromOp>4129struct MustBeDeadConversion : public fir::FIROpConversion<FromOp> {4130 explicit MustBeDeadConversion(const fir::LLVMTypeConverter &lowering,4131 const fir::FIRToLLVMPassOptions &options)4132 : fir::FIROpConversion<FromOp>(lowering, options) {}4133 using OpAdaptor = typename FromOp::Adaptor;4134 4135 llvm::LogicalResult4136 matchAndRewrite(FromOp op, OpAdaptor adaptor,4137 mlir::ConversionPatternRewriter &rewriter) const final {4138 if (!op->getUses().empty())4139 return rewriter.notifyMatchFailure(op, "op must be dead");4140 rewriter.eraseOp(op);4141 return mlir::success();4142 }4143};4144 4145struct ShapeOpConversion : public MustBeDeadConversion<fir::ShapeOp> {4146 using MustBeDeadConversion::MustBeDeadConversion;4147};4148 4149struct ShapeShiftOpConversion : public MustBeDeadConversion<fir::ShapeShiftOp> {4150 using MustBeDeadConversion::MustBeDeadConversion;4151};4152 4153struct ShiftOpConversion : public MustBeDeadConversion<fir::ShiftOp> {4154 using MustBeDeadConversion::MustBeDeadConversion;4155};4156 4157struct SliceOpConversion : public MustBeDeadConversion<fir::SliceOp> {4158 using MustBeDeadConversion::MustBeDeadConversion;4159};4160 4161} // namespace4162 4163namespace {4164class RenameMSVCLibmCallees4165 : public mlir::OpRewritePattern<mlir::LLVM::CallOp> {4166public:4167 using OpRewritePattern::OpRewritePattern;4168 4169 llvm::LogicalResult4170 matchAndRewrite(mlir::LLVM::CallOp op,4171 mlir::PatternRewriter &rewriter) const override {4172 rewriter.startOpModification(op);4173 auto callee = op.getCallee();4174 if (callee)4175 if (*callee == "hypotf")4176 op.setCalleeAttr(mlir::SymbolRefAttr::get(op.getContext(), "_hypotf"));4177 4178 rewriter.finalizeOpModification(op);4179 return mlir::success();4180 }4181};4182 4183class RenameMSVCLibmFuncs4184 : public mlir::OpRewritePattern<mlir::LLVM::LLVMFuncOp> {4185public:4186 using OpRewritePattern::OpRewritePattern;4187 4188 llvm::LogicalResult4189 matchAndRewrite(mlir::LLVM::LLVMFuncOp op,4190 mlir::PatternRewriter &rewriter) const override {4191 rewriter.startOpModification(op);4192 if (op.getSymName() == "hypotf")4193 op.setSymNameAttr(rewriter.getStringAttr("_hypotf"));4194 rewriter.finalizeOpModification(op);4195 return mlir::success();4196 }4197};4198} // namespace4199 4200namespace {4201/// Convert FIR dialect to LLVM dialect4202///4203/// This pass lowers all FIR dialect operations to LLVM IR dialect. An4204/// MLIR pass is used to lower residual Std dialect to LLVM IR dialect.4205class FIRToLLVMLowering4206 : public fir::impl::FIRToLLVMLoweringBase<FIRToLLVMLowering> {4207public:4208 FIRToLLVMLowering() = default;4209 FIRToLLVMLowering(fir::FIRToLLVMPassOptions options) : options{options} {}4210 mlir::ModuleOp getModule() { return getOperation(); }4211 4212 void runOnOperation() override final {4213 auto mod = getModule();4214 if (!forcedTargetTriple.empty())4215 fir::setTargetTriple(mod, forcedTargetTriple);4216 4217 if (!forcedDataLayout.empty()) {4218 llvm::DataLayout dl(forcedDataLayout);4219 fir::support::setMLIRDataLayout(mod, dl);4220 }4221 4222 if (!forcedTargetCPU.empty())4223 fir::setTargetCPU(mod, forcedTargetCPU);4224 4225 if (!forcedTuneCPU.empty())4226 fir::setTuneCPU(mod, forcedTuneCPU);4227 4228 if (!forcedTargetFeatures.empty())4229 fir::setTargetFeatures(mod, forcedTargetFeatures);4230 4231 if (typeDescriptorsRenamedForAssembly)4232 options.typeDescriptorsRenamedForAssembly =4233 typeDescriptorsRenamedForAssembly;4234 4235 // Run dynamic pass pipeline for converting Math dialect4236 // operations into other dialects (llvm, func, etc.).4237 // Some conversions of Math operations cannot be done4238 // by just using conversion patterns. This is true for4239 // conversions that affect the ModuleOp, e.g. create new4240 // function operations in it. We have to run such conversions4241 // as passes here.4242 mlir::OpPassManager mathConversionPM("builtin.module");4243 4244 bool isAMDGCN = fir::getTargetTriple(mod).isAMDGCN();4245 // If compiling for AMD target some math operations must be lowered to AMD4246 // GPU library calls, the rest can be converted to LLVM intrinsics, which4247 // is handled in the mathToLLVM conversion. The lowering to libm calls is4248 // not needed since all math operations are handled this way.4249 if (isAMDGCN) {4250 mathConversionPM.addPass(mlir::createConvertMathToROCDL());4251 mathConversionPM.addPass(mlir::createConvertComplexToROCDLLibraryCalls());4252 }4253 4254 // Convert math::FPowI operations to inline implementation4255 // only if the exponent's width is greater than 32, otherwise,4256 // it will be lowered to LLVM intrinsic operation by a later conversion.4257 mlir::ConvertMathToFuncsOptions mathToFuncsOptions{};4258 mathToFuncsOptions.minWidthOfFPowIExponent = 33;4259 mathConversionPM.addPass(4260 mlir::createConvertMathToFuncs(mathToFuncsOptions));4261 4262 mlir::ConvertComplexToStandardPassOptions complexToStandardOptions{};4263 if (options.ComplexRange ==4264 Fortran::frontend::CodeGenOptions::ComplexRangeKind::CX_Basic) {4265 complexToStandardOptions.complexRange =4266 mlir::complex::ComplexRangeFlags::basic;4267 } else if (options.ComplexRange == Fortran::frontend::CodeGenOptions::4268 ComplexRangeKind::CX_Improved) {4269 complexToStandardOptions.complexRange =4270 mlir::complex::ComplexRangeFlags::improved;4271 }4272 mathConversionPM.addPass(4273 mlir::createConvertComplexToStandardPass(complexToStandardOptions));4274 4275 // Convert Math dialect operations into LLVM dialect operations.4276 // There is no way to prefer MathToLLVM patterns over MathToLibm4277 // patterns (applied below), so we have to run MathToLLVM conversion here.4278 mathConversionPM.addNestedPass<mlir::func::FuncOp>(4279 mlir::createConvertMathToLLVMPass());4280 if (mlir::failed(runPipeline(mathConversionPM, mod)))4281 return signalPassFailure();4282 4283 std::optional<mlir::DataLayout> dl =4284 fir::support::getOrSetMLIRDataLayout(mod, /*allowDefaultLayout=*/true);4285 if (!dl) {4286 mlir::emitError(mod.getLoc(),4287 "module operation must carry a data layout attribute "4288 "to generate llvm IR from FIR");4289 signalPassFailure();4290 return;4291 }4292 4293 auto *context = getModule().getContext();4294 fir::LLVMTypeConverter typeConverter{getModule(),4295 options.applyTBAA || applyTBAA,4296 options.forceUnifiedTBAATree, *dl};4297 mlir::RewritePatternSet pattern(context);4298 fir::populateFIRToLLVMConversionPatterns(typeConverter, pattern, options);4299 mlir::populateFuncToLLVMConversionPatterns(typeConverter, pattern);4300 mlir::populateOpenMPToLLVMConversionPatterns(typeConverter, pattern);4301 mlir::arith::populateArithToLLVMConversionPatterns(typeConverter, pattern);4302 mlir::cf::populateControlFlowToLLVMConversionPatterns(typeConverter,4303 pattern);4304 mlir::cf::populateAssertToLLVMConversionPattern(typeConverter, pattern);4305 // Math operations that have not been converted yet must be converted4306 // to Libm.4307 if (!isAMDGCN)4308 mlir::populateMathToLibmConversionPatterns(pattern);4309 mlir::populateComplexToLLVMConversionPatterns(typeConverter, pattern);4310 mlir::index::populateIndexToLLVMConversionPatterns(typeConverter, pattern);4311 mlir::populateVectorToLLVMConversionPatterns(typeConverter, pattern);4312 4313 // Flang specific overloads for OpenMP operations, to allow for special4314 // handling of things like Box types.4315 fir::populateOpenMPFIRToLLVMConversionPatterns(typeConverter, pattern);4316 4317 mlir::ConversionTarget target{*context};4318 target.addLegalDialect<mlir::LLVM::LLVMDialect>();4319 // The OpenMP dialect is legal for Operations without regions, for those4320 // which contains regions it is legal if the region contains only the4321 // LLVM dialect. Add OpenMP dialect as a legal dialect for conversion and4322 // legalize conversion of OpenMP operations without regions.4323 mlir::configureOpenMPToLLVMConversionLegality(target, typeConverter);4324 target.addLegalDialect<mlir::omp::OpenMPDialect>();4325 target.addLegalDialect<mlir::acc::OpenACCDialect>();4326 target.addLegalDialect<mlir::gpu::GPUDialect>();4327 4328 // required NOPs for applying a full conversion4329 target.addLegalOp<mlir::ModuleOp>();4330 4331 // If we're on Windows, we might need to rename some libm calls.4332 bool isMSVC = fir::getTargetTriple(mod).isOSMSVCRT();4333 if (isMSVC) {4334 pattern.insert<RenameMSVCLibmCallees, RenameMSVCLibmFuncs>(context);4335 4336 target.addDynamicallyLegalOp<mlir::LLVM::CallOp>(4337 [](mlir::LLVM::CallOp op) {4338 auto callee = op.getCallee();4339 if (!callee)4340 return true;4341 return *callee != "hypotf";4342 });4343 target.addDynamicallyLegalOp<mlir::LLVM::LLVMFuncOp>(4344 [](mlir::LLVM::LLVMFuncOp op) {4345 return op.getSymName() != "hypotf";4346 });4347 }4348 4349 // apply the patterns4350 if (mlir::failed(mlir::applyFullConversion(getModule(), target,4351 std::move(pattern)))) {4352 signalPassFailure();4353 }4354 4355 // Run pass to add comdats to functions that have weak linkage on relevant4356 // platforms4357 if (fir::getTargetTriple(mod).supportsCOMDAT()) {4358 mlir::OpPassManager comdatPM("builtin.module");4359 comdatPM.addPass(mlir::LLVM::createLLVMAddComdats());4360 if (mlir::failed(runPipeline(comdatPM, mod)))4361 return signalPassFailure();4362 }4363 }4364 4365private:4366 fir::FIRToLLVMPassOptions options;4367};4368 4369/// Lower from LLVM IR dialect to proper LLVM-IR and dump the module4370struct LLVMIRLoweringPass4371 : public mlir::PassWrapper<LLVMIRLoweringPass,4372 mlir::OperationPass<mlir::ModuleOp>> {4373 MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(LLVMIRLoweringPass)4374 4375 LLVMIRLoweringPass(llvm::raw_ostream &output, fir::LLVMIRLoweringPrinter p)4376 : output{output}, printer{p} {}4377 4378 mlir::ModuleOp getModule() { return getOperation(); }4379 4380 void runOnOperation() override final {4381 auto *ctx = getModule().getContext();4382 auto optName = getModule().getName();4383 llvm::LLVMContext llvmCtx;4384 if (auto llvmModule = mlir::translateModuleToLLVMIR(4385 getModule(), llvmCtx, optName ? *optName : "FIRModule")) {4386 printer(*llvmModule, output);4387 return;4388 }4389 4390 mlir::emitError(mlir::UnknownLoc::get(ctx), "could not emit LLVM-IR\n");4391 signalPassFailure();4392 }4393 4394private:4395 llvm::raw_ostream &output;4396 fir::LLVMIRLoweringPrinter printer;4397};4398 4399} // namespace4400 4401std::unique_ptr<mlir::Pass> fir::createFIRToLLVMPass() {4402 return std::make_unique<FIRToLLVMLowering>();4403}4404 4405std::unique_ptr<mlir::Pass>4406fir::createFIRToLLVMPass(fir::FIRToLLVMPassOptions options) {4407 return std::make_unique<FIRToLLVMLowering>(options);4408}4409 4410std::unique_ptr<mlir::Pass>4411fir::createLLVMDialectToLLVMPass(llvm::raw_ostream &output,4412 fir::LLVMIRLoweringPrinter printer) {4413 return std::make_unique<LLVMIRLoweringPass>(output, printer);4414}4415 4416void fir::populateFIRToLLVMConversionPatterns(4417 const fir::LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns,4418 fir::FIRToLLVMPassOptions &options) {4419 patterns.insert<4420 AbsentOpConversion, AddcOpConversion, AddrOfOpConversion,4421 AllocaOpConversion, AllocMemOpConversion, BoxAddrOpConversion,4422 BoxCharLenOpConversion, BoxDimsOpConversion, BoxEleSizeOpConversion,4423 BoxIsAllocOpConversion, BoxIsArrayOpConversion, BoxIsPtrOpConversion,4424 AssumedSizeExtentOpConversion, IsAssumedSizeExtentOpConversion,4425 BoxOffsetOpConversion, BoxProcHostOpConversion, BoxRankOpConversion,4426 BoxTypeCodeOpConversion, BoxTypeDescOpConversion, CallOpConversion,4427 CmpcOpConversion, VolatileCastOpConversion, ConvertOpConversion,4428 CoordinateOpConversion, CopyOpConversion, DTEntryOpConversion,4429 DeclareOpConversion,4430 DoConcurrentSpecifierOpConversion<fir::LocalitySpecifierOp>,4431 DoConcurrentSpecifierOpConversion<fir::DeclareReductionOp>,4432 DivcOpConversion, EmboxOpConversion, EmboxCharOpConversion,4433 EmboxProcOpConversion, ExtractValueOpConversion, FieldIndexOpConversion,4434 FirEndOpConversion, FreeMemOpConversion, GlobalLenOpConversion,4435 GlobalOpConversion, InsertOnRangeOpConversion, IsPresentOpConversion,4436 LenParamIndexOpConversion, LoadOpConversion, MulcOpConversion,4437 NegcOpConversion, NoReassocOpConversion, SelectCaseOpConversion,4438 SelectOpConversion, SelectRankOpConversion, SelectTypeOpConversion,4439 ShapeOpConversion, ShapeShiftOpConversion, ShiftOpConversion,4440 SliceOpConversion, StoreOpConversion, StringLitOpConversion,4441 SubcOpConversion, TypeDescOpConversion, TypeInfoOpConversion,4442 UnboxCharOpConversion, UnboxProcOpConversion, UndefOpConversion,4443 UnreachableOpConversion, XArrayCoorOpConversion, XEmboxOpConversion,4444 XReboxOpConversion, ZeroOpConversion>(converter, options);4445 4446 // Patterns that are populated without a type converter do not trigger4447 // target materializations for the operands of the root op.4448 patterns.insert<HasValueOpConversion, InsertValueOpConversion>(4449 patterns.getContext());4450}4451