Lines Matching full:mlir
9 // Coding style: https://mlir.llvm.org/getting_started/DeveloperGuide/
31 #include "mlir/Conversion/ArithCommon/AttrToLLVMConverter.h"
32 #include "mlir/Conversion/ArithToLLVM/ArithToLLVM.h"
33 #include "mlir/Conversion/ComplexToLLVM/ComplexToLLVM.h"
34 #include "mlir/Conversion/ComplexToStandard/ComplexToStandard.h"
35 #include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h"
36 #include "mlir/Conversion/FuncToLLVM/ConvertFuncToLLVM.h"
37 #include "mlir/Conversion/LLVMCommon/Pattern.h"
38 #include "mlir/Conversion/MathToFuncs/MathToFuncs.h"
39 #include "mlir/Conversion/MathToLLVM/MathToLLVM.h"
40 #include "mlir/Conversion/MathToLibm/MathToLibm.h"
41 #include "mlir/Conversion/MathToROCDL/MathToROCDL.h"
42 #include "mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h"
43 #include "mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h"
44 #include "mlir/Dialect/Arith/IR/Arith.h"
45 #include "mlir/Dialect/DLTI/DLTI.h"
46 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
47 #include "mlir/Dialect/LLVMIR/LLVMAttrs.h"
48 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
49 #include "mlir/Dialect/LLVMIR/Transforms/AddComdats.h"
50 #include "mlir/Dialect/OpenACC/OpenACC.h"
51 #include "mlir/Dialect/OpenMP/OpenMPDialect.h"
52 #include "mlir/IR/BuiltinTypes.h"
53 #include "mlir/IR/Matchers.h"
54 #include "mlir/Pass/Pass.h"
55 #include "mlir/Pass/PassManager.h"
56 #include "mlir/Target/LLVMIR/Import.h"
57 #include "mlir/Target/LLVMIR/ModuleTranslation.h"
76 static inline mlir::Type getLlvmPtrType(mlir::MLIRContext *context,
78 return mlir::LLVM::LLVMPointerType::get(context, addressSpace);
81 static inline mlir::Type getI8Type(mlir::MLIRContext *context) {
82 return mlir::IntegerType::get(context, 8);
85 static mlir::LLVM::ConstantOp
86 genConstantIndex(mlir::Location loc, mlir::Type ity,
87 mlir::ConversionPatternRewriter &rewriter,
90 return rewriter.create<mlir::LLVM::ConstantOp>(loc, ity, cattr);
93 static mlir::Block *createBlock(mlir::ConversionPatternRewriter &rewriter,
94 mlir::Block *insertBefore) {
97 mlir::Region::iterator(insertBefore));
102 static int64_t getConstantIntValue(mlir::Value val) {
108 static unsigned getTypeDescFieldId(mlir::Type ty) {
109 auto isArray = mlir::isa<fir::SequenceType>(fir::dyn_cast_ptrOrBoxEleTy(ty));
112 static unsigned getLenParamFieldId(mlir::Type ty) {
116 static llvm::SmallVector<mlir::NamedAttribute>
117 addLLVMOpBundleAttrs(mlir::ConversionPatternRewriter &rewriter,
118 llvm::ArrayRef<mlir::NamedAttribute> attrs,
120 llvm::SmallVector<mlir::NamedAttribute> newAttrs;
123 for (mlir::NamedAttribute attr : attrs) {
143 mlir::ConversionPatternRewriter &rewriter) const override {
145 rewriter.replaceOpWithNewOp<mlir::LLVM::AddressOfOp>(
147 return mlir::success();
155 static mlir::LLVM::LLVMFuncOp
157 mlir::ConversionPatternRewriter &rewriter) {
158 auto module = op->getParentOfType<mlir::ModuleOp>();
160 if (auto memSizeFunc = module.lookupSymbol<mlir::LLVM::LLVMFuncOp>(name))
170 static mlir::Value
171 genAllocationScaleSize(OP op, mlir::Type ity,
172 mlir::ConversionPatternRewriter &rewriter) {
173 mlir::Location loc = op.getLoc();
174 mlir::Type dataTy = op.getInType();
175 auto seqTy = mlir::dyn_cast<fir::SequenceType>(dataTy);
191 mlir::Value constVal{
204 mlir::ConversionPatternRewriter &rewriter) const override {
206 if (auto fusedLoc = mlir::dyn_cast<mlir::FusedLoc>(declareOp.getLoc())) {
208 mlir::dyn_cast_or_null<mlir::LLVM::DILocalVariableAttr>(
210 rewriter.create<mlir::LLVM::DbgDeclareOp>(memRef.getLoc(), memRef,
215 return mlir::success();
227 mlir::ConversionPatternRewriter &rewriter) const override {
228 mlir::ValueRange operands = adaptor.getOperands();
230 mlir::Type ity = lowerTy().indexType();
232 mlir::Value size = genConstantIndex(loc, ity, rewriter, 1).getResult();
233 mlir::Type firObjType = fir::unwrapRefType(alloc.getType());
234 mlir::Type llvmObjectType = convertObjectType(firObjType);
237 llvm::SmallVector<mlir::Value> lenParams;
240 mlir::Type scalarType = fir::unwrapSequenceType(alloc.getInType());
241 if (auto chrTy = mlir::dyn_cast<fir::CharacterType>(scalarType)) {
247 } else if (auto recTy = mlir::dyn_cast<fir::RecordType>(scalarType)) {
248 mlir::LLVM::LLVMFuncOp memSizeFn =
252 mlir::NamedAttribute attr = rewriter.getNamedAttr(
253 "callee", mlir::SymbolRefAttr::get(memSizeFn));
254 auto call = rewriter.create<mlir::LLVM::CallOp>(
266 rewriter.createOrFold<mlir::LLVM::MulOp>(loc, ity, size, scaleSize);
270 size = rewriter.createOrFold<mlir::LLVM::MulOp>(
278 if (mlir::isa<mlir::LLVM::ConstantOp>(size.getDefiningOp())) {
280 mlir::Operation *parentOp = rewriter.getInsertionBlock()->getParentOp();
281 mlir::Region *parentRegion = rewriter.getInsertionBlock()->getParent();
282 mlir::Block *insertBlock =
289 mlir::Operation *clonedSize = rewriter.clone(*size.getDefiningOp());
299 auto llvmAlloc = rewriter.create<mlir::LLVM::AllocaOp>(
316 rewriter.replaceOpWithNewOp<mlir::LLVM::AddrSpaceCastOp>(
319 return mlir::success();
332 mlir::ConversionPatternRewriter &rewriter) const override {
333 mlir::Value a = adaptor.getOperands()[0];
336 mlir::dyn_cast<fir::BaseBoxType>(boxaddr.getVal().getType())) {
341 rewriter.replaceOpWithNewOp<mlir::LLVM::ExtractValueOp>(boxaddr, a, 0);
343 return mlir::success();
354 mlir::ConversionPatternRewriter &rewriter) const override {
355 mlir::Value boxChar = adaptor.getOperands()[0];
356 mlir::Location loc = boxChar.getLoc();
357 mlir::Type returnValTy = boxCharLen.getResult().getType();
360 auto len = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, boxChar,
362 mlir::Value lenAfterCast = integerCast(loc, rewriter, returnValTy, len);
365 return mlir::success();
377 mlir::ConversionPatternRewriter &rewriter) const override {
378 llvm::SmallVector<mlir::Type, 3> resultTypes = {
388 return mlir::success();
399 mlir::ConversionPatternRewriter &rewriter) const override {
400 mlir::Value box = adaptor.getOperands()[0];
406 return mlir::success();
417 mlir::ConversionPatternRewriter &rewriter) const override {
418 mlir::Value box = adaptor.getOperands()[0];
421 mlir::Value check =
424 return mlir::success();
435 mlir::ConversionPatternRewriter &rewriter) const override {
436 mlir::Value a = adaptor.getOperands()[0];
439 mlir::Value rank = getRankFromBox(loc, boxTyPair, a, rewriter);
440 mlir::Value c0 = genConstantIndex(loc, rank.getType(), rewriter, 0);
441 rewriter.replaceOpWithNewOp<mlir::LLVM::ICmpOp>(
442 boxisarray, mlir::LLVM::ICmpPredicate::ne, rank, c0);
443 return mlir::success();
454 mlir::ConversionPatternRewriter &rewriter) const override {
455 mlir::Value box = adaptor.getOperands()[0];
458 mlir::Value check =
461 return mlir::success();
472 mlir::ConversionPatternRewriter &rewriter) const override {
473 mlir::Value a = adaptor.getOperands()[0];
475 mlir::Type ty = convertType(boxrank.getType());
478 mlir::Value rank = getRankFromBox(loc, boxTyPair, a, rewriter);
479 mlir::Value result = integerCast(loc, rewriter, ty, rank);
481 return mlir::success();
494 mlir::ConversionPatternRewriter &rewriter) const override {
496 return mlir::failure();
508 mlir::ConversionPatternRewriter &rewriter) const override {
509 mlir::Value box = adaptor.getOperands()[0];
514 return mlir::success();
526 mlir::ConversionPatternRewriter &rewriter) const override {
527 mlir::Value box = adaptor.getOperands()[0];
534 return mlir::success();
544 mlir::ConversionPatternRewriter &rewriter) const override {
547 if (mlir::isa<mlir::StringAttr>(attr)) {
548 rewriter.replaceOpWithNewOp<mlir::LLVM::ConstantOp>(constop, ty, attr);
549 return mlir::success();
552 auto charTy = mlir::cast<fir::CharacterType>(constop.getType());
554 mlir::Type intTy = rewriter.getIntegerType(bits);
555 mlir::Location loc = constop.getLoc();
556 mlir::Value cst = rewriter.create<mlir::LLVM::UndefOp>(loc, ty);
557 if (auto arr = mlir::dyn_cast<mlir::DenseElementsAttr>(attr)) {
558 cst = rewriter.create<mlir::LLVM::ConstantOp>(loc, ty, arr);
559 } else if (auto arr = mlir::dyn_cast<mlir::ArrayAttr>(attr)) {
562 auto elemAttr = mlir::IntegerAttr::get(
564 mlir::cast<mlir::IntegerAttr>(a.value()).getValue().zextOrTrunc(
567 rewriter.create<mlir::LLVM::ConstantOp>(loc, intTy, elemAttr);
568 cst = rewriter.create<mlir::LLVM::InsertValueOp>(loc, cst, elemCst,
572 return mlir::failure();
575 return mlir::success();
585 mlir::ConversionPatternRewriter &rewriter) const override {
586 llvm::SmallVector<mlir::Type> resultTys;
590 mlir::arith::AttrConvertFastMathToLLVM<fir::CallOp, mlir::LLVM::CallOp>
592 rewriter.replaceOpWithNewOp<mlir::LLVM::CallOp>(
596 return mlir::success();
601 static mlir::Type getComplexEleTy(mlir::Type complex) {
602 return mlir::cast<mlir::ComplexType>(complex).getElementType();
616 mlir::ConversionPatternRewriter &rewriter) const override {
617 mlir::ValueRange operands = adaptor.getOperands();
618 mlir::Type resTy = convertType(cmp.getType());
619 mlir::Location loc = cmp.getLoc();
620 mlir::LLVM::FastmathFlags fmf =
621 mlir::arith::convertArithFastMathFlagsToLLVM(cmp.getFastmath());
622 mlir::LLVM::FCmpPredicate pred =
623 static_cast<mlir::LLVM::FCmpPredicate>(cmp.getPredicate());
624 auto rcp = rewriter.create<mlir::LLVM::FCmpOp>(
626 rewriter.create<mlir::LLVM::ExtractValueOp>(loc, operands[0], 0),
627 rewriter.create<mlir::LLVM::ExtractValueOp>(loc, operands[1], 0), fmf);
628 auto icp = rewriter.create<mlir::LLVM::FCmpOp>(
630 rewriter.create<mlir::LLVM::ExtractValueOp>(loc, operands[0], 1),
631 rewriter.create<mlir::LLVM::ExtractValueOp>(loc, operands[1], 1), fmf);
632 llvm::SmallVector<mlir::Value, 2> cp = {rcp, icp};
634 case mlir::arith::CmpFPredicate::OEQ: // .EQ.
635 rewriter.replaceOpWithNewOp<mlir::LLVM::AndOp>(cmp, resTy, cp);
637 case mlir::arith::CmpFPredicate::UNE: // .NE.
638 rewriter.replaceOpWithNewOp<mlir::LLVM::OrOp>(cmp, resTy, cp);
644 return mlir::success();
652 static bool isFloatingPointTy(mlir::Type ty) {
653 return mlir::isa<mlir::FloatType>(ty);
658 mlir::ConversionPatternRewriter &rewriter) const override {
663 mlir::Value op0 = adaptor.getOperands()[0];
667 return mlir::success();
671 auto i1Type = mlir::IntegerType::get(convert.getContext(), 1);
673 if (mlir::isa<fir::RecordType>(toFirTy)) {
677 assert(mlir::cast<fir::RecordType>(fromFirTy).getTypeList() ==
678 mlir::cast<fir::RecordType>(toFirTy).getTypeList() &&
681 auto toStTy = mlir::cast<mlir::LLVM::LLVMStructType>(toTy);
682 mlir::Value val = rewriter.create<mlir::LLVM::UndefOp>(loc, toStTy);
687 int64_t index = mlir::cast<mlir::IntegerAttr>(attr).getInt();
689 rewriter.create<mlir::LLVM::ExtractValueOp>(loc, op0, index);
691 rewriter.create<mlir::LLVM::InsertValueOp>(loc, val, extVal, index);
695 return mlir::success();
698 if (mlir::isa<fir::LogicalType>(fromFirTy) ||
699 mlir::isa<fir::LogicalType>(toFirTy)) {
712 if (!mlir::isa<mlir::IntegerType>(fromTy) ||
713 !mlir::isa<mlir::IntegerType>(toTy))
714 return mlir::emitError(loc)
720 mlir::Value normVal =
723 return mlir::success();
729 rewriter.replaceOpWithNewOp<mlir::LLVM::ZExtOp>(convert, toTy, op0);
730 return mlir::success();
734 mlir::Value zero = genConstantIndex(loc, fromTy, rewriter, 0);
735 auto isTrue = rewriter.create<mlir::LLVM::ICmpOp>(
736 loc, mlir::LLVM::ICmpPredicate::ne, op0, zero);
741 rewriter.replaceOpWithNewOp<mlir::LLVM::ZExtOp>(convert, toTy, isTrue);
745 return mlir::success();
750 return mlir::success();
752 auto convertFpToFp = [&](mlir::Value val, unsigned fromBits,
753 unsigned toBits, mlir::Type toTy) -> mlir::Value {
757 mlir::emitError(loc,
763 return rewriter.create<mlir::LLVM::FPTruncOp>(loc, toTy, val);
764 return rewriter.create<mlir::LLVM::FPExtOp>(loc, toTy, val);
771 auto rp = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, op0, 0);
772 auto ip = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, op0, 1);
774 auto fromBits = mlir::LLVM::getPrimitiveTypeSizeInBits(ty);
775 auto toBits = mlir::LLVM::getPrimitiveTypeSizeInBits(nt);
778 auto un = rewriter.create<mlir::LLVM::UndefOp>(loc, toTy);
779 auto i1 = rewriter.create<mlir::LLVM::InsertValueOp>(loc, un, rc, 0);
780 rewriter.replaceOpWithNewOp<mlir::LLVM::InsertValueOp>(convert, i1, ic,
782 return mlir::success();
788 auto fromBits = mlir::LLVM::getPrimitiveTypeSizeInBits(fromTy);
789 auto toBits = mlir::LLVM::getPrimitiveTypeSizeInBits(toTy);
792 return mlir::success();
794 if (mlir::isa<mlir::IntegerType>(toTy)) {
796 rewriter.replaceOpWithNewOp<mlir::LLVM::FPToUIOp>(convert, toTy, op0);
798 rewriter.replaceOpWithNewOp<mlir::LLVM::FPToSIOp>(convert, toTy, op0);
799 return mlir::success();
801 } else if (mlir::isa<mlir::IntegerType>(fromTy)) {
803 if (mlir::isa<mlir::IntegerType>(toTy)) {
804 auto fromBits = mlir::LLVM::getPrimitiveTypeSizeInBits(fromTy);
805 auto toBits = mlir::LLVM::getPrimitiveTypeSizeInBits(toTy);
808 rewriter.replaceOpWithNewOp<mlir::LLVM::TruncOp>(convert, toTy, op0);
809 return mlir::success();
812 rewriter.replaceOpWithNewOp<mlir::LLVM::ZExtOp>(convert, toTy, op0);
813 return mlir::success();
815 rewriter.replaceOpWithNewOp<mlir::LLVM::SExtOp>(convert, toTy, op0);
816 return mlir::success();
821 rewriter.replaceOpWithNewOp<mlir::LLVM::UIToFPOp>(convert, toTy, op0);
823 rewriter.replaceOpWithNewOp<mlir::LLVM::SIToFPOp>(convert, toTy, op0);
824 return mlir::success();
827 if (mlir::isa<mlir::LLVM::LLVMPointerType>(toTy)) {
828 rewriter.replaceOpWithNewOp<mlir::LLVM::IntToPtrOp>(convert, toTy, op0);
829 return mlir::success();
831 } else if (mlir::isa<mlir::LLVM::LLVMPointerType>(fromTy)) {
833 if (mlir::isa<mlir::IntegerType>(toTy)) {
834 rewriter.replaceOpWithNewOp<mlir::LLVM::PtrToIntOp>(convert, toTy, op0);
835 return mlir::success();
838 if (mlir::isa<mlir::LLVM::LLVMPointerType>(toTy)) {
839 rewriter.replaceOpWithNewOp<mlir::LLVM::BitcastOp>(convert, toTy, op0);
840 return mlir::success();
856 mlir::ConversionPatternRewriter &rewriter) const override {
858 return mlir::success();
869 mlir::ConversionPatternRewriter &rewriter) const override {
871 return mlir::success();
881 mlir::ConversionPatternRewriter &rewriter) const override {
883 return mlir::failure();
895 mlir::ConversionPatternRewriter &rewriter) const override {
909 mlir::ConversionPatternRewriter &rewriter) const override {
910 mlir::ValueRange operands = adaptor.getOperands();
912 mlir::Value charBuffer = operands[0];
913 mlir::Value charBufferLen = operands[1];
915 mlir::Location loc = emboxChar.getLoc();
916 mlir::Type llvmStructTy = convertType(emboxChar.getType());
917 auto llvmStruct = rewriter.create<mlir::LLVM::UndefOp>(loc, llvmStructTy);
919 mlir::Type lenTy =
920 mlir::cast<mlir::LLVM::LLVMStructType>(llvmStructTy).getBody()[1];
921 mlir::Value lenAfterCast = integerCast(loc, rewriter, lenTy, charBufferLen);
923 mlir::Type addrTy =
924 mlir::cast<mlir::LLVM::LLVMStructType>(llvmStructTy).getBody()[0];
927 rewriter.create<mlir::LLVM::BitcastOp>(loc, addrTy, charBuffer);
929 auto insertBufferOp = rewriter.create<mlir::LLVM::InsertValueOp>(
931 rewriter.replaceOpWithNewOp<mlir::LLVM::InsertValueOp>(
934 return mlir::success();
940 static mlir::SymbolRefAttr
942 mlir::ConversionPatternRewriter &rewriter) {
945 mod.template lookupSymbol<mlir::LLVM::LLVMFuncOp>(mallocName))
946 return mlir::SymbolRefAttr::get(mallocFunc);
948 mod.template lookupSymbol<mlir::func::FuncOp>(mallocName))
949 return mlir::SymbolRefAttr::get(userMalloc);
951 mlir::OpBuilder moduleBuilder(mod.getBodyRegion());
952 auto indexType = mlir::IntegerType::get(op.getContext(), 64);
953 auto mallocDecl = moduleBuilder.create<mlir::LLVM::LLVMFuncOp>(
955 mlir::LLVM::LLVMFunctionType::get(getLlvmPtrType(op.getContext()),
958 return mlir::SymbolRefAttr::get(mallocDecl);
962 static mlir::SymbolRefAttr
963 getMalloc(fir::AllocMemOp op, mlir::ConversionPatternRewriter &rewriter) {
964 if (auto mod = op->getParentOfType<mlir::gpu::GPUModuleOp>())
966 auto mod = op->getParentOfType<mlir::ModuleOp>();
974 static mlir::Value
975 computeElementDistance(mlir::Location loc, mlir::Type llvmObjectType,
976 mlir::Type idxTy,
977 mlir::ConversionPatternRewriter &rewriter) {
979 // mlir::LLVM::getPrimitiveTypeSizeInBits() for the element type here. For
980 // example, it returns 10 bytes for mlir::Float80Type for targets where it
982 // mlir::DataLayout::getTypeABIAlignment(), but DataLayout is not being set
988 auto nullPtr = rewriter.create<mlir::LLVM::ZeroOp>(loc, llvmPtrTy);
989 auto gep = rewriter.create<mlir::LLVM::GEPOp>(
991 llvm::ArrayRef<mlir::LLVM::GEPArg>{1});
992 return rewriter.create<mlir::LLVM::PtrToIntOp>(loc, idxTy, gep);
998 static mlir::Value
999 genTypeStrideInBytes(mlir::Location loc, mlir::Type idxTy,
1000 mlir::ConversionPatternRewriter &rewriter,
1001 mlir::Type llTy) {
1013 mlir::ConversionPatternRewriter &rewriter) const override {
1014 mlir::Type heapTy = heap.getType();
1015 mlir::Location loc = heap.getLoc();
1017 mlir::Type dataTy = fir::unwrapRefType(heapTy);
1018 mlir::Type llvmObjectTy = convertObjectType(dataTy);
1021 mlir::Value size = genTypeSizeInBytes(loc, ity, rewriter, llvmObjectTy);
1023 size = rewriter.create<mlir::LLVM::MulOp>(loc, ity, size, scaleSize);
1024 for (mlir::Value opnd : adaptor.getOperands())
1025 size = rewriter.create<mlir::LLVM::MulOp>(
1028 rewriter.replaceOpWithNewOp<mlir::LLVM::CallOp>(
1031 return mlir::success();
1037 mlir::Value genTypeSizeInBytes(mlir::Location loc, mlir::Type idxTy,
1038 mlir::ConversionPatternRewriter &rewriter,
1039 mlir::Type llTy) const {
1047 static mlir::SymbolRefAttr
1049 mlir::ConversionPatternRewriter &rewriter) {
1053 mod.template lookupSymbol<mlir::LLVM::LLVMFuncOp>(freeName))
1054 return mlir::SymbolRefAttr::get(freeFunc);
1056 mod.template lookupSymbol<mlir::func::FuncOp>(freeName))
1057 return mlir::SymbolRefAttr::get(freeDefinedByUser);
1059 mlir::OpBuilder moduleBuilder(mod.getBodyRegion());
1060 auto voidType = mlir::LLVM::LLVMVoidType::get(op.getContext());
1061 auto freeDecl = moduleBuilder.create<mlir::LLVM::LLVMFuncOp>(
1063 mlir::LLVM::LLVMFunctionType::get(voidType,
1066 return mlir::SymbolRefAttr::get(freeDecl);
1069 static mlir::SymbolRefAttr getFree(fir::FreeMemOp op,
1070 mlir::ConversionPatternRewriter &rewriter) {
1071 if (auto mod = op->getParentOfType<mlir::gpu::GPUModuleOp>())
1073 auto mod = op->getParentOfType<mlir::ModuleOp>();
1077 static unsigned getDimension(mlir::LLVM::LLVMArrayType ty) {
1080 mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(ty.getElementType());
1081 eleTy; eleTy = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(
1094 mlir::ConversionPatternRewriter &rewriter) const override {
1095 mlir::Location loc = freemem.getLoc();
1097 rewriter.create<mlir::LLVM::CallOp>(
1098 loc, mlir::TypeRange{}, mlir::ValueRange{adaptor.getHeapref()},
1101 return mlir::success();
1107 static llvm::SmallVector<mlir::Value>
1108 convertSubcomponentIndices(mlir::Location loc, mlir::Type eleTy,
1109 mlir::ValueRange indices,
1110 mlir::Type *retTy = nullptr) {
1111 llvm::SmallVector<mlir::Value> result;
1112 llvm::SmallVector<mlir::Value> arrayIndices;
1122 for (mlir::Value index : indices) {
1125 if (auto structTy = mlir::dyn_cast<mlir::LLVM::LLVMStructType>(eleTy)) {
1133 mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(eleTy)) {
1145 static mlir::Value genSourceFile(mlir::Location loc, mlir::ModuleOp mod,
1146 mlir::ConversionPatternRewriter &rewriter) {
1147 auto ptrTy = mlir::LLVM::LLVMPointerType::get(rewriter.getContext());
1148 if (auto flc = mlir::dyn_cast<mlir::FileLineColLoc>(loc)) {
1153 return rewriter.create<mlir::LLVM::AddressOfOp>(loc, ptrTy, g.getName());
1154 } else if (auto g = mod.lookupSymbol<mlir::LLVM::GlobalOp>(globalName)) {
1155 return rewriter.create<mlir::LLVM::AddressOfOp>(loc, ptrTy, g.getName());
1160 auto arrayTy = mlir::LLVM::LLVMArrayType::get(
1161 mlir::IntegerType::get(rewriter.getContext(), 8), fn.size());
1162 mlir::LLVM::GlobalOp globalOp = rewriter.create<mlir::LLVM::GlobalOp>(
1163 loc, arrayTy, /*constant=*/true, mlir::LLVM::Linkage::Linkonce,
1164 globalName, mlir::Attribute());
1166 mlir::Region ®ion = globalOp.getInitializerRegion();
1167 mlir::Block *block = rewriter.createBlock(®ion);
1169 mlir::Value constValue = rewriter.create<mlir::LLVM::ConstantOp>(
1171 rewriter.create<mlir::LLVM::ReturnOp>(loc, constValue);
1173 return rewriter.create<mlir::LLVM::AddressOfOp>(loc, ptrTy,
1176 return rewriter.create<mlir::LLVM::ZeroOp>(loc, ptrTy);
1179 static mlir::Value genSourceLine(mlir::Location loc,
1180 mlir::ConversionPatternRewriter &rewriter) {
1181 if (auto flc = mlir::dyn_cast<mlir::FileLineColLoc>(loc))
1182 return rewriter.create<mlir::LLVM::ConstantOp>(loc, rewriter.getI32Type(),
1184 return rewriter.create<mlir::LLVM::ConstantOp>(loc, rewriter.getI32Type(), 0);
1187 static mlir::Value
1188 genCUFAllocDescriptor(mlir::Location loc,
1189 mlir::ConversionPatternRewriter &rewriter,
1190 mlir::ModuleOp mod, fir::BaseBoxType boxTy,
1192 std::optional<mlir::DataLayout> dl =
1195 mlir::emitError(mod.getLoc(),
1199 mlir::Value sourceFile = genSourceFile(loc, mod, rewriter);
1200 mlir::Value sourceLine = genSourceLine(loc, rewriter);
1202 mlir::MLIRContext *ctx = mod.getContext();
1204 mlir::LLVM::LLVMPointerType llvmPointerType =
1205 mlir::LLVM::LLVMPointerType::get(ctx);
1206 mlir::Type llvmInt32Type = mlir::IntegerType::get(ctx, 32);
1207 mlir::Type llvmIntPtrType =
1208 mlir::IntegerType::get(ctx, typeConverter.getPointerBitwidth(0));
1209 auto fctTy = mlir::LLVM::LLVMFunctionType::get(
1212 auto llvmFunc = mod.lookupSymbol<mlir::LLVM::LLVMFuncOp>(
1215 mod.lookupSymbol<mlir::func::FuncOp>(RTNAME_STRING(CUFAllocDescriptor));
1217 mlir::OpBuilder::atBlockEnd(mod.getBody())
1218 .create<mlir::LLVM::LLVMFuncOp>(loc, RTNAME_STRING(CUFAllocDescriptor),
1221 mlir::Type structTy = typeConverter.convertBoxTypeAsStruct(boxTy);
1223 mlir::Value sizeInBytes =
1227 .create<mlir::LLVM::CallOp>(loc, fctTy, RTNAME_STRING(CUFAllocDescriptor),
1240 if (mlir::isa<fir::PointerType>(eleTy))
1242 if (mlir::isa<fir::HeapType>(eleTy))
1247 mlir::Value getCharacterByteSize(mlir::Location loc,
1248 mlir::ConversionPatternRewriter &rewriter,
1250 mlir::ValueRange lenParams) const {
1251 auto i64Ty = mlir::IntegerType::get(rewriter.getContext(), 64);
1252 mlir::Value size =
1260 return rewriter.create<mlir::LLVM::MulOp>(loc, i64Ty, size, len64);
1264 std::tuple<mlir::Value, mlir::Value> getSizeAndTypeCode(
1265 mlir::Location loc, mlir::ConversionPatternRewriter &rewriter,
1266 mlir::Type boxEleTy, mlir::ValueRange lenParams = {}) const {
1267 auto i64Ty = mlir::IntegerType::get(rewriter.getContext(), 64);
1270 if (auto seqTy = mlir::dyn_cast<fir::SequenceType>(boxEleTy))
1272 if (mlir::isa<mlir::NoneType>(
1274 return {rewriter.create<mlir::LLVM::ConstantOp>(loc, i64Ty, 0),
1276 mlir::Value typeCodeVal = this->genConstantOffset(
1280 mlir::dyn_cast<fir::LogicalType>(boxEleTy) || fir::isa_real(boxEleTy) ||
1285 if (auto charTy = mlir::dyn_cast<fir::CharacterType>(boxEleTy))
1292 if (mlir::isa<fir::RecordType>(boxEleTy))
1300 mlir::Value insertField(mlir::ConversionPatternRewriter &rewriter,
1301 mlir::Location loc, mlir::Value dest,
1303 mlir::Value value, bool bitcast = false) const {
1309 return rewriter.create<mlir::LLVM::InsertValueOp>(loc, dest, value,
1313 inline mlir::Value
1314 insertBaseAddress(mlir::ConversionPatternRewriter &rewriter,
1315 mlir::Location loc, mlir::Value dest,
1316 mlir::Value base) const {
1321 inline mlir::Value insertLowerBound(mlir::ConversionPatternRewriter &rewriter,
1322 mlir::Location loc, mlir::Value dest,
1323 unsigned dim, mlir::Value lb) const {
1328 inline mlir::Value insertExtent(mlir::ConversionPatternRewriter &rewriter,
1329 mlir::Location loc, mlir::Value dest,
1330 unsigned dim, mlir::Value extent) const {
1335 inline mlir::Value insertStride(mlir::ConversionPatternRewriter &rewriter,
1336 mlir::Location loc, mlir::Value dest,
1337 unsigned dim, mlir::Value stride) const {
1345 mlir::Value
1346 getTypeDescriptor(ModOpTy mod, mlir::ConversionPatternRewriter &rewriter,
1347 mlir::Location loc, fir::RecordType recType) const {
1352 mlir::Type llvmPtrTy = ::getLlvmPtrType(mod.getContext());
1354 return rewriter.create<mlir::LLVM::AddressOfOp>(loc, llvmPtrTy,
1357 if (auto global = mod.template lookupSymbol<mlir::LLVM::GlobalOp>(name)) {
1359 return rewriter.create<mlir::LLVM::AddressOfOp>(loc, llvmPtrTy,
1369 return rewriter.create<mlir::LLVM::ZeroOp>(loc, llvmPtrTy);
1373 mlir::Value populateDescriptor(mlir::Location loc, ModOpTy mod,
1374 fir::BaseBoxType boxTy, mlir::Type inputType,
1375 mlir::ConversionPatternRewriter &rewriter,
1376 unsigned rank, mlir::Value eleSize,
1377 mlir::Value cfiTy, mlir::Value typeDesc,
1379 mlir::Value extraField = {}) const {
1383 mlir::Value descriptor =
1384 rewriter.create<mlir::LLVM::UndefOp>(loc, llvmBoxTy);
1402 auto maskAttr = mlir::IntegerAttr::get(
1405 mlir::LLVM::ConstantOp mask = rewriter.create<mlir::LLVM::ConstantOp>(
1407 extraField = rewriter.create<mlir::LLVM::OrOp>(loc, extraField, mask);
1409 auto maskAttr = mlir::IntegerAttr::get(
1412 mlir::LLVM::ConstantOp mask = rewriter.create<mlir::LLVM::ConstantOp>(
1414 extraField = rewriter.create<mlir::LLVM::AndOp>(loc, extraField, mask);
1433 mlir::Type innerType = fir::unwrapInnerType(inputType);
1434 if (innerType && mlir::isa<fir::RecordType>(innerType)) {
1435 auto recTy = mlir::dyn_cast<fir::RecordType>(innerType);
1440 typeDesc = rewriter.create<mlir::LLVM::ZeroOp>(
1456 mlir::Value zero =
1466 std::tuple<fir::BaseBoxType, mlir::Value, mlir::Value>
1467 consDescriptorPrefix(BOX box, mlir::Type inputType,
1468 mlir::ConversionPatternRewriter &rewriter, unsigned rank,
1469 [[maybe_unused]] mlir::ValueRange substrParams,
1470 mlir::ValueRange lenParams, mlir::Value sourceBox = {},
1471 mlir::Type sourceBoxType = {}) const {
1473 auto boxTy = mlir::dyn_cast<fir::BaseBoxType>(box.getType());
1476 llvm::SmallVector<mlir::Value> typeparams = lenParams;
1496 mlir::Value typeDesc;
1497 mlir::Value extraField;
1504 mlir::Type idxTy = this->lowerTy().indexType();
1513 mlir::Value descriptor;
1514 if (auto gpuMod = box->template getParentOfType<mlir::gpu::GPUModuleOp>())
1518 else if (auto mod = box->template getParentOfType<mlir::ModuleOp>())
1526 std::tuple<fir::BaseBoxType, mlir::Value, mlir::Value>
1527 consDescriptorPrefix(fir::cg::XReboxOp box, mlir::Value loweredBox,
1528 mlir::ConversionPatternRewriter &rewriter, unsigned rank,
1529 mlir::ValueRange substrParams,
1530 mlir::ValueRange lenParams,
1531 mlir::Value typeDesc = {}) const {
1533 auto boxTy = mlir::dyn_cast<fir::BaseBoxType>(box.getType());
1534 auto inputBoxTy = mlir::dyn_cast<fir::BaseBoxType>(box.getBox().getType());
1536 llvm::SmallVector<mlir::Value> typeparams = lenParams;
1548 mlir::Type idxTy = this->lowerTy().indexType();
1561 mlir::Value extraField =
1564 mlir::Value descriptor;
1565 if (auto gpuMod = box->template getParentOfType<mlir::gpu::GPUModuleOp>())
1570 else if (auto mod = box->template getParentOfType<mlir::ModuleOp>())
1583 mlir::Value
1584 genBoxOffsetGep(mlir::ConversionPatternRewriter &rewriter, mlir::Location loc,
1585 mlir::Value base, mlir::Type llvmBaseObjectType,
1586 mlir::Value outerOffset, mlir::ValueRange cstInteriorIndices,
1587 mlir::ValueRange componentIndices,
1588 std::optional<mlir::Value> substringOffset) const {
1589 llvm::SmallVector<mlir::LLVM::GEPArg> gepArgs{outerOffset};
1590 mlir::Type resultTy = llvmBaseObjectType;
1592 for (mlir::Value interiorIndex : llvm::reverse(cstInteriorIndices)) {
1593 auto arrayTy = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(resultTy);
1601 llvm::SmallVector<mlir::Value> gepIndices =
1605 if (auto arrayTy = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(resultTy)) {
1616 mlir::Type outterOffsetTy =
1617 llvm::cast<mlir::Value>(gepArgs[0]).getType();
1618 mlir::Value cast =
1621 gepArgs[0] = rewriter.create<mlir::LLVM::AddOp>(
1622 loc, outterOffsetTy, llvm::cast<mlir::Value>(gepArgs[0]), cast);
1625 mlir::Type llvmPtrTy = ::getLlvmPtrType(resultTy.getContext());
1626 return rewriter.create<mlir::LLVM::GEPOp>(
1632 getSubcomponentIndices(BOX xbox, mlir::Value memref,
1633 mlir::ValueRange operands,
1634 mlir::SmallVectorImpl<mlir::Value> &indices) const {
1647 static bool isInGlobalOp(mlir::ConversionPatternRewriter &rewriter) {
1650 mlir::isa<mlir::LLVM::GlobalOp>(thisBlock->getParentOp());
1656 mlir::Value
1657 placeInMemoryIfNotGlobalInit(mlir::ConversionPatternRewriter &rewriter,
1658 mlir::Location loc, mlir::Type boxTy,
1659 mlir::Value boxValue,
1663 mlir::Type llvmBoxTy = boxValue.getType();
1664 mlir::Value storage;
1666 auto mod = boxValue.getDefiningOp()->getParentOfType<mlir::ModuleOp>();
1667 auto baseBoxTy = mlir::dyn_cast<fir::BaseBoxType>(boxTy);
1674 auto storeOp = rewriter.create<mlir::LLVM::StoreOp>(loc, boxValue, storage);
1680 mlir::Value computeTripletExtent(mlir::ConversionPatternRewriter &rewriter,
1681 mlir::Location loc, mlir::Value lb,
1682 mlir::Value ub, mlir::Value step,
1683 mlir::Value zero, mlir::Type type) const {
1688 mlir::Value extent = rewriter.create<mlir::LLVM::SubOp>(loc, type, ub, lb);
1689 extent = rewriter.create<mlir::LLVM::AddOp>(loc, type, extent, step);
1690 extent = rewriter.create<mlir::LLVM::SDivOp>(loc, type, extent, step);
1693 auto cmp = rewriter.create<mlir::LLVM::ICmpOp>(
1694 loc, mlir::LLVM::ICmpPredicate::sgt, extent, zero);
1695 return rewriter.create<mlir::LLVM::SelectOp>(loc, cmp, extent, zero);
1706 mlir::ConversionPatternRewriter &rewriter) const override {
1707 mlir::ValueRange operands = adaptor.getOperands();
1708 mlir::Value sourceBox;
1709 mlir::Type sourceBoxType;
1717 /*rank=*/0, /*substrParams=*/mlir::ValueRange{},
1723 return mlir::failure();
1728 return mlir::success();
1732 static bool isDeviceAllocation(mlir::Value val, mlir::Value adaptorVal) {
1733 if (auto loadOp = mlir::dyn_cast_or_null<fir::LoadOp>(val.getDefiningOp()))
1736 mlir::dyn_cast_or_null<fir::BoxAddrOp>(val.getDefiningOp()))
1739 mlir::dyn_cast_or_null<fir::ConvertOp>(val.getDefiningOp()))
1742 if (auto blockArg = llvm::cast<mlir::BlockArgument>(adaptorVal)) {
1745 if (auto func = mlir::dyn_cast_or_null<mlir::FunctionOpInterface>(
1751 mlir::dyn_cast<cuf::DataAttributeAttr>(attr.getValue());
1761 if (auto callOp = mlir::dyn_cast_or_null<fir::CallOp>(val.getDefiningOp()))
1777 mlir::ConversionPatternRewriter &rewriter) const override {
1778 mlir::ValueRange operands = adaptor.getOperands();
1779 mlir::Value sourceBox;
1780 mlir::Type sourceBoxType;
1790 auto i64Ty = mlir::IntegerType::get(xbox.getContext(), 64);
1797 mlir::Location loc = xbox.getLoc();
1798 mlir::Value zero = genConstantIndex(loc, i64Ty, rewriter, 0);
1799 mlir::Value one = genConstantIndex(loc, i64Ty, rewriter, 1);
1800 mlir::Value prevPtrOff = one;
1801 mlir::Type eleTy = boxTy.getEleTy();
1803 llvm::SmallVector<mlir::Value> cstInteriorIndices;
1805 mlir::Value ptrOffset = zero;
1806 mlir::Type memEleTy = fir::dyn_cast_ptrEleTy(xbox.getMemref().getType());
1807 assert(mlir::isa<fir::SequenceType>(memEleTy));
1808 auto seqTy = mlir::cast<fir::SequenceType>(memEleTy);
1809 mlir::Type seqEleTy = seqTy.getEleTy();
1812 if (auto charTy = mlir::dyn_cast<fir::CharacterType>(seqEleTy)) {
1818 } else if (mlir::isa<fir::RecordType>(seqEleTy)) {
1834 mlir::Value prevDimByteStride = resultEleSize;
1843 auto charTy = mlir::cast<fir::CharacterType>(seqEleTy);
1858 mlir::Value extent =
1860 mlir::Value outerExtent = extent;
1863 mlir::Value off =
1865 mlir::Value adj = one;
1868 auto ao = rewriter.create<mlir::LLVM::SubOp>(loc, i64Ty, off, adj);
1873 rewriter.create<mlir::LLVM::MulOp>(loc, i64Ty, ao, prevPtrOff);
1875 rewriter.create<mlir::LLVM::AddOp>(loc, i64Ty, dimOff, ptrOffset);
1877 if (mlir::isa_and_nonnull<fir::UndefOp>(
1892 mlir::Value lb = zero;
1894 mlir::isa<fir::PointerType, fir::HeapType>(eleTy);
1905 auto extentIsEmpty = rewriter.create<mlir::LLVM::ICmpOp>(
1906 loc, mlir::LLVM::ICmpPredicate::eq, extent, zero);
1907 lb = rewriter.create<mlir::LLVM::SelectOp>(loc, extentIsEmpty, one,
1915 mlir::Value step = prevDimByteStride;
1917 mlir::Value sliceStep =
1920 rewriter.create<mlir::LLVM::MulOp>(loc, i64Ty, step, sliceStep);
1927 prevDimByteStride = rewriter.create<mlir::LLVM::MulOp>(
1930 prevPtrOff = rewriter.create<mlir::LLVM::MulOp>(loc, i64Ty, prevPtrOff,
1942 mlir::Value base = adaptor.getMemref();
1945 llvm::SmallVector<mlir::Value> fieldIndices;
1946 std::optional<mlir::Value> substringOffset;
1951 mlir::Type llvmBaseType =
1959 mlir::Value result = placeInMemoryIfNotGlobalInit(
1963 return mlir::success();
1980 mlir::ConversionPatternRewriter &rewriter) const override {
1981 mlir::Location loc = rebox.getLoc();
1982 mlir::Type idxTy = lowerTy().indexType();
1983 mlir::Value loweredBox = adaptor.getOperands()[0];
1984 mlir::ValueRange operands = adaptor.getOperands();
1990 // fir.box was translated to an llvm.ptr<llvm.struct<>> and the MLIR pass
1995 loweredBox.getDefiningOp<mlir::UnrealizedConversionCastOp>())
2001 llvm::SmallVector<mlir::Value, 2> lenParams;
2002 mlir::Type inputEleTy = getInputEleTy(rebox);
2003 if (auto charTy = mlir::dyn_cast<fir::CharacterType>(inputEleTy)) {
2005 mlir::Value len =
2009 mlir::Value len = getElementSizeFromBox(loc, idxTy, inputBoxTyPair,
2014 mlir::Value width =
2016 len = rewriter.create<mlir::LLVM::SDivOp>(loc, idxTy, len, width);
2020 } else if (auto recTy = mlir::dyn_cast<fir::RecordType>(inputEleTy)) {
2026 mlir::Value typeDescAddr;
2027 if (mlir::isa<fir::ClassType>(inputBoxTyPair.fir) &&
2028 mlir::isa<fir::ClassType>(rebox.getType()))
2037 llvm::SmallVector<mlir::Value> inputExtents;
2038 llvm::SmallVector<mlir::Value> inputStrides;
2041 llvm::SmallVector<mlir::Value, 3> dimInfo =
2048 mlir::Value baseAddr =
2063 mlir::Type destBoxTy, mlir::Value dest, mlir::Value base,
2064 mlir::ValueRange lbounds, mlir::ValueRange extents,
2065 mlir::ValueRange strides,
2066 mlir::ConversionPatternRewriter &rewriter) const {
2067 mlir::Location loc = rebox.getLoc();
2068 mlir::Value zero =
2070 mlir::Value one = genConstantIndex(loc, lowerTy().indexType(), rewriter, 1);
2072 mlir::Value extent = std::get<0>(iter.value());
2074 mlir::Value lb = one;
2077 auto extentIsEmpty = rewriter.create<mlir::LLVM::ICmpOp>(
2078 loc, mlir::LLVM::ICmpPredicate::eq, extent, zero);
2079 lb = rewriter.create<mlir::LLVM::SelectOp>(loc, extentIsEmpty, one, lb);
2086 mlir::Value result = placeInMemoryIfNotGlobalInit(
2090 return mlir::success();
2095 sliceBox(fir::cg::XReboxOp rebox, OpAdaptor adaptor, mlir::Type destBoxTy,
2096 mlir::Value dest, mlir::Value base, mlir::ValueRange inputExtents,
2097 mlir::ValueRange inputStrides, mlir::ValueRange operands,
2098 mlir::ConversionPatternRewriter &rewriter) const {
2099 mlir::Location loc = rebox.getLoc();
2100 mlir::Type byteTy = ::getI8Type(rebox.getContext());
2101 mlir::Type idxTy = lowerTy().indexType();
2102 mlir::Value zero = genConstantIndex(loc, idxTy, rewriter, 0);
2106 mlir::Type inputEleTy = getInputEleTy(rebox);
2107 mlir::Type llvmBaseObjectType = convertType(inputEleTy);
2108 llvm::SmallVector<mlir::Value> fieldIndices;
2109 std::optional<mlir::Value> substringOffset;
2128 llvm::SmallVector<mlir::Value> slicedExtents;
2129 llvm::SmallVector<mlir::Value> slicedStrides;
2130 mlir::Value one = genConstantIndex(loc, idxTy, rewriter, 1);
2138 mlir::Value sliceLb =
2140 mlir::Value inputStride = *strideOps; // already idxTy
2142 mlir::Value sliceOrigin =
2146 mlir::Value diff =
2147 rewriter.create<mlir::LLVM::SubOp>(loc, idxTy, sliceLb, sliceOrigin);
2148 mlir::Value offset =
2149 rewriter.create<mlir::LLVM::MulOp>(loc, idxTy, diff, inputStride);
2154 mlir::Value upper = operands[sliceOps + 1];
2156 !mlir::isa_and_nonnull<mlir::LLVM::UndefOp>(upper.getDefiningOp());
2158 mlir::Value step =
2161 mlir::Value sliceUb = integerCast(loc, rewriter, idxTy, upper);
2162 mlir::Value extent = computeTripletExtent(rewriter, loc, sliceLb,
2166 mlir::Value stride =
2167 rewriter.create<mlir::LLVM::MulOp>(loc, idxTy, step, inputStride);
2179 reshapeBox(fir::cg::XReboxOp rebox, OpAdaptor adaptor, mlir::Type destBoxTy,
2180 mlir::Value dest, mlir::Value base, mlir::ValueRange inputExtents,
2181 mlir::ValueRange inputStrides, mlir::ValueRange operands,
2182 mlir::ConversionPatternRewriter &rewriter) const {
2183 mlir::ValueRange reboxShifts{
2193 mlir::Location loc = rebox.getLoc();
2195 llvm::SmallVector<mlir::Value> newStrides;
2196 llvm::SmallVector<mlir::Value> newExtents;
2197 mlir::Type idxTy = lowerTy().indexType();
2202 mlir::Value stride = inputStrides.empty()
2206 mlir::Value rawExtent = operands[rebox.getShapeOperandIndex() + i];
2207 mlir::Value extent = integerCast(loc, rewriter, idxTy, rawExtent);
2211 stride = rewriter.create<mlir::LLVM::MulOp>(loc, idxTy, extent, stride);
2218 static mlir::Type getInputEleTy(fir::cg::XReboxOp rebox) {
2220 if (auto seqTy = mlir::dyn_cast<fir::SequenceType>(ty))
2233 mlir::ConversionPatternRewriter &rewriter) const override {
2235 return mlir::failure();
2244 mlir::Type ty) {
2248 if (auto seq = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(ty)) {
2256 } else if (auto st = mlir::dyn_cast<mlir::LLVM::LLVMStructType>(ty)) {
2265 collectIndices(mlir::ConversionPatternRewriter &rewriter,
2266 mlir::ArrayAttr arrAttr) {
2269 if (auto intAttr = mlir::dyn_cast<mlir::IntegerAttr>(*i)) {
2272 auto fieldName = mlir::cast<mlir::StringAttr>(*i).getValue();
2274 auto ty = mlir::cast<mlir::TypeAttr>(*i).getValue();
2275 auto index = mlir::cast<fir::RecordType>(ty).getFieldIndex(fieldName);
2283 static mlir::Type getArrayElementType(mlir::LLVM::LLVMArrayType ty) {
2285 while (auto arrTy = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(eleTy))
2299 doRewrite(fir::ExtractValueOp extractVal, mlir::Type ty, OpAdaptor adaptor,
2300 mlir::ConversionPatternRewriter &rewriter) const override {
2301 mlir::ValueRange operands = adaptor.getOperands();
2304 rewriter.replaceOpWithNewOp<mlir::LLVM::ExtractValueOp>(
2306 return mlir::success();
2313 : public mlir::OpConversionPattern<fir::InsertValueOp>,
2319 mlir::ConversionPatternRewriter &rewriter) const override {
2320 mlir::ValueRange operands = adaptor.getOperands();
2323 rewriter.replaceOpWithNewOp<mlir::LLVM::InsertValueOp>(
2325 return mlir::success();
2346 doRewrite(fir::InsertOnRangeOp range, mlir::Type ty, OpAdaptor adaptor,
2347 mlir::ConversionPatternRewriter &rewriter) const override {
2353 while (auto t = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(type)) {
2362 mlir::DenseIntElementsAttr coor = range.getCoor();
2371 mlir::Value lastOp = adaptor.getOperands()[0];
2372 mlir::Value insertVal = adaptor.getOperands()[1];
2375 lastOp = rewriter.create<mlir::LLVM::InsertValueOp>(
2381 rewriter.replaceOpWithNewOp<mlir::LLVM::InsertValueOp>(
2384 return mlir::success();
2399 doRewrite(fir::cg::XArrayCoorOp coor, mlir::Type llvmPtrTy, OpAdaptor adaptor,
2400 mlir::ConversionPatternRewriter &rewriter) const override {
2402 mlir::ValueRange operands = adaptor.getOperands();
2408 mlir::Type idxTy = lowerTy().indexType();
2414 mlir::Value one = genConstantIndex(loc, idxTy, rewriter, 1);
2415 mlir::Value prevExt = one;
2416 mlir::Value offset = genConstantIndex(loc, idxTy, rewriter, 0);
2420 mlir::isa<fir::BaseBoxType>(coor.getMemref().getType());
2423 mlir::LLVM::IntegerOverflowFlags nsw =
2424 mlir::LLVM::IntegerOverflowFlags::nsw;
2429 mlir::Value index =
2431 mlir::Value lb =
2434 mlir::Value step = one;
2439 mlir::Value originalUb = *(sliceOps + 1);
2441 !mlir::isa_and_nonnull<fir::UndefOp>(originalUb.getDefiningOp());
2445 auto idx = rewriter.create<mlir::LLVM::SubOp>(loc, idxTy, index, lb, nsw);
2446 mlir::Value diff =
2447 rewriter.create<mlir::LLVM::MulOp>(loc, idxTy, idx, step, nsw);
2449 mlir::Value sliceLb =
2452 rewriter.create<mlir::LLVM::SubOp>(loc, idxTy, sliceLb, lb, nsw);
2453 diff = rewriter.create<mlir::LLVM::AddOp>(loc, idxTy, diff, adj, nsw);
2459 mlir::Value stride =
2462 rewriter.create<mlir::LLVM::MulOp>(loc, idxTy, diff, stride, nsw);
2464 rewriter.create<mlir::LLVM::AddOp>(loc, idxTy, sc, offset, nsw);
2468 rewriter.create<mlir::LLVM::MulOp>(loc, idxTy, diff, prevExt, nsw);
2470 rewriter.create<mlir::LLVM::AddOp>(loc, idxTy, sc, offset, nsw);
2474 prevExt = rewriter.create<mlir::LLVM::MulOp>(loc, idxTy, prevExt,
2483 mlir::Type byteTy = ::getI8Type(coor.getContext());
2484 mlir::Value base =
2486 llvm::SmallVector<mlir::LLVM::GEPArg> args{offset};
2487 auto addr = rewriter.create<mlir::LLVM::GEPOp>(loc, llvmPtrTy, byteTy,
2491 return mlir::success();
2496 mlir::Type elementType =
2499 mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(elementType))
2508 llvm::SmallVector<mlir::Value> indices = convertSubcomponentIndices(
2513 rewriter.replaceOpWithNewOp<mlir::LLVM::GEPOp>(coor, llvmPtrTy,
2515 return mlir::success();
2521 mlir::Type objectTy = fir::unwrapRefType(coor.getMemref().getType());
2522 mlir::Type eleType = fir::unwrapSequenceType(objectTy);
2523 mlir::Type gepObjectType = convertType(eleType);
2524 llvm::SmallVector<mlir::LLVM::GEPArg> args;
2535 offset = rewriter.create<mlir::LLVM::MulOp>(loc, idxTy, offset,
2545 llvm::SmallVector<mlir::Value> indices = convertSubcomponentIndices(
2551 rewriter.replaceOpWithNewOp<mlir::LLVM::GEPOp>(
2553 return mlir::success();
2568 doRewrite(fir::CoordinateOp coor, mlir::Type ty, OpAdaptor adaptor,
2569 mlir::ConversionPatternRewriter &rewriter) const override {
2570 mlir::ValueRange operands = adaptor.getOperands();
2572 mlir::Location loc = coor.getLoc();
2573 mlir::Value base = operands[0];
2574 mlir::Type baseObjectTy = coor.getBaseType();
2575 mlir::Type objectTy = fir::dyn_cast_ptrOrBoxEleTy(baseObjectTy);
2577 mlir::Type llvmObjectTy = convertType(objectTy);
2582 mlir::Value gep =
2585 return mlir::success();
2589 if (mlir::dyn_cast<fir::BaseBoxType>(baseObjectTy))
2593 if (mlir::isa<fir::ReferenceType, fir::PointerType, fir::HeapType>(
2601 static unsigned getFieldNumber(fir::RecordType ty, mlir::Value op) {
2604 ->getAttrOfType<mlir::IntegerAttr>("field")
2609 static bool hasSubDimensions(mlir::Type type) {
2610 return mlir::isa<fir::SequenceType, fir::RecordType, mlir::TupleType>(type);
2618 static bool supportedCoordinate(mlir::Type type, mlir::ValueRange coors) {
2624 mlir::Value nxtOpnd = coors[i];
2625 if (auto arrTy = mlir::dyn_cast<fir::SequenceType>(type)) {
2629 } else if (auto recTy = mlir::dyn_cast<fir::RecordType>(type)) {
2632 } else if (auto tupTy = mlir::dyn_cast<mlir::TupleType>(type)) {
2647 static bool arraysHaveKnownShape(mlir::Type type, mlir::ValueRange coors) {
2649 mlir::Value nxtOpnd = coors[i];
2650 if (auto arrTy = mlir::dyn_cast<fir::SequenceType>(type)) {
2655 } else if (auto strTy = mlir::dyn_cast<fir::RecordType>(type)) {
2657 } else if (auto strTy = mlir::dyn_cast<mlir::TupleType>(type)) {
2668 doRewriteBox(fir::CoordinateOp coor, mlir::ValueRange operands,
2669 mlir::Location loc,
2670 mlir::ConversionPatternRewriter &rewriter) const {
2671 mlir::Type boxObjTy = coor.getBaseType();
2672 assert(mlir::dyn_cast<fir::BaseBoxType>(boxObjTy) &&
2676 mlir::Value boxBaseAddr = operands[0];
2683 mlir::Operation *coordinateDef =
2685 if (mlir::isa_and_nonnull<fir::LenParamIndexOp>(coordinateDef))
2705 mlir::Value resultAddr =
2709 mlir::Type llvmPtrTy = ::getLlvmPtrType(coor.getContext());
2710 mlir::Type byteTy = ::getI8Type(coor.getContext());
2711 mlir::LLVM::IntegerOverflowFlags nsw =
2712 mlir::LLVM::IntegerOverflowFlags::nsw;
2715 if (auto arrTy = mlir::dyn_cast<fir::SequenceType>(cpnTy)) {
2723 mlir::Value off = genConstantIndex(loc, idxTy, rewriter, 0);
2726 mlir::Value stride = getStrideFromBox(loc, boxTyPair, operands[0],
2728 auto sc = rewriter.create<mlir::LLVM::MulOp>(
2730 off = rewriter.create<mlir::LLVM::AddOp>(loc, idxTy, sc, off, nsw);
2732 resultAddr = rewriter.create<mlir::LLVM::GEPOp>(
2734 llvm::ArrayRef<mlir::LLVM::GEPArg>{off});
2737 } else if (auto recTy = mlir::dyn_cast<fir::RecordType>(cpnTy)) {
2738 mlir::Value nxtOpnd = operands[i];
2741 resultAddr = rewriter.create<mlir::LLVM::GEPOp>(
2743 llvm::ArrayRef<mlir::LLVM::GEPArg>{0, nxtOpnd});
2750 return mlir::success();
2754 doRewriteRefOrPtr(fir::CoordinateOp coor, mlir::Type llvmObjectTy,
2755 mlir::ValueRange operands, mlir::Location loc,
2756 mlir::ConversionPatternRewriter &rewriter) const {
2757 mlir::Type baseObjectTy = coor.getBaseType();
2760 mlir::Type cpnTy = fir::dyn_cast_ptrOrBoxEleTy(baseObjectTy);
2772 if (auto arrTy = mlir::dyn_cast<fir::SequenceType>(cpnTy)) {
2792 return mlir::emitError(
2796 llvm::SmallVector<mlir::LLVM::GEPArg> offs;
2801 llvm::SmallVector<mlir::Value> arrIdx;
2803 mlir::Value nxtOpnd = operands[i];
2806 return mlir::emitError(loc, "invalid coordinate/check failed");
2816 cpnTy = mlir::cast<fir::SequenceType>(cpnTy).getElementType();
2823 if (auto arrTy = mlir::dyn_cast<fir::SequenceType>(cpnTy)) {
2830 cpnTy = mlir::cast<fir::SequenceType>(cpnTy).getElementType();
2836 if (auto recTy = mlir::dyn_cast<fir::RecordType>(cpnTy))
2838 else if (auto tupTy = mlir::dyn_cast<mlir::TupleType>(cpnTy))
2847 mlir::Value base = operands[0];
2848 mlir::Value retval = genGEP(loc, llvmObjectTy, rewriter, base, offs);
2850 return mlir::success();
2853 return mlir::emitError(
2866 mlir::ConversionPatternRewriter &rewriter) const override {
2867 auto recTy = mlir::cast<fir::RecordType>(field.getOnType());
2873 rewriter.replaceOp(field, mlir::ValueRange{genConstantOffset(
2875 return mlir::success();
2881 mlir::FlatSymbolRefAttr symAttr = mlir::SymbolRefAttr::get(
2883 mlir::NamedAttribute callAttr = rewriter.getNamedAttr("callee", symAttr);
2884 mlir::NamedAttribute fieldAttr = rewriter.getNamedAttr(
2885 "field", mlir::IntegerAttr::get(lowerTy().indexType(), index));
2886 rewriter.replaceOpWithNewOp<mlir::LLVM::CallOp>(
2890 return mlir::success();
2907 mlir::ConversionPatternRewriter &rewriter) const override {
2909 return mlir::failure();
2919 mlir::ConversionPatternRewriter &rewriter) const override {
2920 mlir::Type inTy = typeDescOp.getInType();
2921 assert(mlir::isa<fir::RecordType>(inTy) && "expecting fir.type");
2922 auto recordType = mlir::dyn_cast<fir::RecordType>(inTy);
2923 auto module = typeDescOp.getOperation()->getParentOfType<mlir::ModuleOp>();
2930 if (auto global = module.lookupSymbol<mlir::LLVM::GlobalOp>(typeDescName)) {
2931 rewriter.replaceOpWithNewOp<mlir::LLVM::AddressOfOp>(
2933 return mlir::success();
2935 rewriter.replaceOpWithNewOp<mlir::LLVM::AddressOfOp>(
2937 return mlir::success();
2939 return mlir::failure();
2945 : public mlir::OpConversionPattern<fir::HasValueOp> {
2950 mlir::ConversionPatternRewriter &rewriter) const override {
2951 rewriter.replaceOpWithNewOp<mlir::LLVM::ReturnOp>(op,
2953 return mlir::success();
2968 static inline bool attributeTypeIsCompatible(mlir::MLIRContext *ctx,
2969 mlir::Attribute attr,
2970 mlir::Type ty) {
2974 auto intOrFpEleAttr = mlir::dyn_cast<mlir::DenseIntOrFPElementsAttr>(attr);
2977 auto tensorTy = mlir::dyn_cast<mlir::TensorType>(intOrFpEleAttr.getType());
2980 mlir::Type attrEleTy =
2981 mlir::LLVMTypeConverter(ctx).convertType(tensorTy.getElementType());
2984 auto arrTy = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(ty);
2987 mlir::Type eleTy = arrTy.getElementType();
2988 while ((arrTy = mlir::dyn_cast<mlir::LLVM::LLVMArrayType>(eleTy)))
3003 mlir::ConversionPatternRewriter &rewriter) const override {
3005 llvm::SmallVector<mlir::Attribute> dbgExprs;
3007 if (auto fusedLoc = mlir::dyn_cast<mlir::FusedLoc>(global.getLoc())) {
3008 if (auto gvExprAttr = mlir::dyn_cast_if_present<mlir::ArrayAttr>(
3010 for (auto attr : gvExprAttr.getAsRange<mlir::Attribute>())
3012 mlir::dyn_cast<mlir::LLVM::DIGlobalVariableExpressionAttr>(
3019 if (auto boxType = mlir::dyn_cast<fir::BaseBoxType>(global.getType()))
3022 mlir::Attribute initAttr = global.getInitVal().value_or(mlir::Attribute());
3026 mlir::SymbolRefAttr comdat;
3027 llvm::ArrayRef<mlir::NamedAttribute> attrs;
3028 auto g = rewriter.create<mlir::LLVM::GlobalOp>(
3035 auto module = global->getParentOfType<mlir::ModuleOp>();
3036 auto gpuMod = global->getParentOfType<mlir::gpu::GPUModuleOp>();
3039 (linkage == mlir::LLVM::Linkage::Linkonce ||
3040 linkage == mlir::LLVM::Linkage::LinkonceODR) &&
3064 auto constant = mlir::dyn_cast<mlir::arith::ConstantOp>(op);
3066 auto convertOp = mlir::dyn_cast<fir::ConvertOp>(op);
3069 constant = mlir::cast<mlir::arith::ConstantOp>(
3072 mlir::Type vecType = mlir::VectorType::get(
3074 auto denseAttr = mlir::DenseElementsAttr::get(
3075 mlir::cast<mlir::ShapedType>(vecType), constant.getValue());
3077 rewriter.replaceOpWithNewOp<mlir::arith::ConstantOp>(
3083 return mlir::success();
3086 bool isFullRange(mlir::DenseIntElementsAttr indexes,
3103 mlir::LLVM::Linkage
3108 return mlir::LLVM::Linkage::Internal;
3110 return mlir::LLVM::Linkage::Linkonce;
3112 return mlir::LLVM::Linkage::LinkonceODR;
3114 return mlir::LLVM::Linkage::Common;
3116 return mlir::LLVM::Linkage::Weak;
3118 return mlir::LLVM::Linkage::External;
3122 static void addComdat(mlir::LLVM::GlobalOp &global,
3123 mlir::ConversionPatternRewriter &rewriter,
3124 mlir::ModuleOp module) {
3126 mlir::LLVM::ComdatOp comdatOp =
3127 module.lookupSymbol<mlir::LLVM::ComdatOp>(comdatName);
3130 rewriter.create<mlir::LLVM::ComdatOp>(module.getLoc(), comdatName);
3132 if (auto select = comdatOp.lookupSymbol<mlir::LLVM::ComdatSelectorOp>(
3135 mlir::OpBuilder::InsertionGuard guard(rewriter);
3137 auto selectorOp = rewriter.create<mlir::LLVM::ComdatSelectorOp>(
3139 mlir::LLVM::comdat::Comdat::Any);
3140 global.setComdatAttr(mlir::SymbolRefAttr::get(
3142 mlir::FlatSymbolRefAttr::get(selectorOp.getSymNameAttr())));
3152 mlir::ConversionPatternRewriter &rewriter) const override {
3154 mlir::Type llvmLoadTy = convertObjectType(load.getType());
3155 if (auto boxTy = mlir::dyn_cast<fir::BaseBoxType>(load.getType())) {
3162 mlir::Value newBoxStorage;
3163 mlir::Location loc = load.getLoc();
3164 if (auto callOp = mlir::dyn_cast_or_null<mlir::LLVM::CallOp>(
3171 auto mod = load->getParentOfType<mlir::ModuleOp>();
3181 mlir::Value boxSize =
3183 auto memcpy = rewriter.create<mlir::LLVM::MemcpyOp>(
3186 if (std::optional<mlir::ArrayAttr> optionalTag = load.getTbaa())
3192 auto loadOp = rewriter.create<mlir::LLVM::LoadOp>(
3194 if (std::optional<mlir::ArrayAttr> optionalTag = load.getTbaa())
3200 return mlir::success();
3212 mlir::ConversionPatternRewriter &rewriter) const override {
3214 return mlir::success();
3218 static void genCondBrOp(mlir::Location loc, mlir::Value cmp, mlir::Block *dest,
3219 std::optional<mlir::ValueRange> destOps,
3220 mlir::ConversionPatternRewriter &rewriter,
3221 mlir::Block *newBlock) {
3223 rewriter.create<mlir::LLVM::CondBrOp>(loc, cmp, dest, *destOps, newBlock,
3224 mlir::ValueRange());
3226 rewriter.create<mlir::LLVM::CondBrOp>(loc, cmp, dest, newBlock);
3230 static void genBrOp(A caseOp, mlir::Block *dest, std::optional<B> destOps,
3231 mlir::ConversionPatternRewriter &rewriter) {
3233 rewriter.replaceOpWithNewOp<mlir::LLVM::BrOp>(caseOp, *destOps, dest);
3235 rewriter.replaceOpWithNewOp<mlir::LLVM::BrOp>(caseOp, std::nullopt, dest);
3238 static void genCaseLadderStep(mlir::Location loc, mlir::Value cmp,
3239 mlir::Block *dest,
3240 std::optional<mlir::ValueRange> destOps,
3241 mlir::ConversionPatternRewriter &rewriter) {
3273 mlir::ConversionPatternRewriter &rewriter) const override {
3275 llvm::ArrayRef<mlir::Attribute> cases = caseOp.getCases().getValue();
3278 if (mlir::isa<fir::CharacterType>(ty)) {
3280 return mlir::failure();
3282 mlir::Value selector = caseOp.getSelector(adaptor.getOperands());
3285 mlir::Block *dest = caseOp.getSuccessor(t);
3286 std::optional<mlir::ValueRange> destOps =
3288 std::optional<mlir::ValueRange> cmpOps =
3290 mlir::Attribute attr = cases[t];
3291 assert(mlir::isa<mlir::UnitAttr>(attr) || cmpOps.has_value());
3292 if (mlir::isa<fir::PointIntervalAttr>(attr)) {
3293 auto cmp = rewriter.create<mlir::LLVM::ICmpOp>(
3294 loc, mlir::LLVM::ICmpPredicate::eq, selector, cmpOps->front());
3298 if (mlir::isa<fir::LowerBoundAttr>(attr)) {
3299 auto cmp = rewriter.create<mlir::LLVM::ICmpOp>(
3300 loc, mlir::LLVM::ICmpPredicate::sle, cmpOps->front(), selector);
3304 if (mlir::isa<fir::UpperBoundAttr>(attr)) {
3305 auto cmp = rewriter.create<mlir::LLVM::ICmpOp>(
3306 loc, mlir::LLVM::ICmpPredicate::sle, selector, cmpOps->front());
3310 if (mlir::isa<fir::ClosedIntervalAttr>(attr)) {
3311 mlir::Value caseArg0 = *cmpOps->begin();
3312 auto cmp0 = rewriter.create<mlir::LLVM::ICmpOp>(
3313 loc, mlir::LLVM::ICmpPredicate::sle, caseArg0, selector);
3318 rewriter.create<mlir::LLVM::CondBrOp>(loc, cmp0, newBlock1, newBlock2);
3320 mlir::Value caseArg1 = *(cmpOps->begin() + 1);
3321 auto cmp1 = rewriter.create<mlir::LLVM::ICmpOp>(
3322 loc, mlir::LLVM::ICmpPredicate::sle, selector, caseArg1);
3327 assert(mlir::isa<mlir::UnitAttr>(attr));
3331 return mlir::success();
3338 static llvm::FailureOr<mlir::Block *>
3339 getConvertedBlock(mlir::ConversionPatternRewriter &rewriter,
3340 const mlir::TypeConverter *converter,
3341 mlir::Operation *branchOp, mlir::Block *block,
3342 mlir::TypeRange expectedTypes) {
3351 std::optional<mlir::TypeConverter::SignatureConversion> conversion =
3367 mlir::ConversionPatternRewriter &rewriter,
3368 const mlir::TypeConverter *converter) {
3371 mlir::Value selector = adaptor.getSelector();
3375 llvm::SmallVector<mlir::Block *> destinations;
3376 llvm::SmallVector<mlir::ValueRange> destinationsOperands;
3377 mlir::Block *defaultDestination;
3378 mlir::ValueRange defaultOperands;
3382 mlir::Block *dest = select.getSuccessor(t);
3384 const mlir::Attribute &attr = cases[t];
3385 if (auto intAttr = mlir::dyn_cast<mlir::IntegerAttr>(attr)) {
3386 destinationsOperands.push_back(destOps ? *destOps : mlir::ValueRange{});
3389 mlir::TypeRange(destinationsOperands.back()));
3390 if (mlir::failed(convertedBlock))
3391 return mlir::failure();
3396 assert(mlir::dyn_cast_or_null<mlir::UnitAttr>(attr));
3398 defaultOperands = destOps ? *destOps : mlir::ValueRange{};
3400 mlir::TypeRange(defaultOperands));
3401 if (mlir::failed(convertedBlock))
3402 return mlir::failure();
3408 selector = rewriter.create<mlir::LLVM::TruncOp>(loc, rewriter.getI32Type(),
3411 rewriter.replaceOpWithNewOp<mlir::LLVM::SwitchOp>(
3419 return mlir::success();
3428 mlir::ConversionPatternRewriter &rewriter) const override {
3440 mlir::ConversionPatternRewriter &rewriter) const override {
3452 mlir::ConversionPatternRewriter &rewriter) const override {
3453 mlir::emitError(select.getLoc(),
3455 return mlir::failure();
3465 mlir::ConversionPatternRewriter &rewriter) const override {
3466 mlir::Location loc = store.getLoc();
3467 mlir::Type storeTy = store.getValue().getType();
3468 mlir::Value llvmValue = adaptor.getValue();
3469 mlir::Value llvmMemref = adaptor.getMemref();
3470 mlir::LLVM::AliasAnalysisOpInterface newOp;
3471 if (auto boxTy = mlir::dyn_cast<fir::BaseBoxType>(storeTy)) {
3472 mlir::Type llvmBoxTy = lowerTy().convertBoxTypeAsStruct(boxTy);
3476 mlir::Value boxSize =
3478 newOp = rewriter.create<mlir::LLVM::MemcpyOp>(
3481 newOp = rewriter.create<mlir::LLVM::StoreOp>(loc, llvmValue, llvmMemref);
3483 if (std::optional<mlir::ArrayAttr> optionalTag = store.getTbaa())
3488 return mlir::success();
3501 mlir::ConversionPatternRewriter &rewriter) const override {
3502 mlir::Type lenTy = convertType(unboxchar.getType(1));
3503 mlir::Value tuple = adaptor.getOperands()[0];
3505 mlir::Location loc = unboxchar.getLoc();
3506 mlir::Value ptrToBuffer =
3507 rewriter.create<mlir::LLVM::ExtractValueOp>(loc, tuple, 0);
3509 auto len = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, tuple, 1);
3510 mlir::Value lenAfterCast = integerCast(loc, rewriter, lenTy, len);
3513 llvm::ArrayRef<mlir::Value>{ptrToBuffer, lenAfterCast});
3514 return mlir::success();
3526 mlir::ConversionPatternRewriter &rewriter) const override {
3528 return mlir::failure();
3538 mlir::ConversionPatternRewriter &rewriter) const override {
3539 rewriter.replaceOpWithNewOp<mlir::LLVM::UndefOp>(
3541 return mlir::success();
3550 mlir::ConversionPatternRewriter &rewriter) const override {
3551 mlir::Type ty = convertType(zero.getType());
3552 rewriter.replaceOpWithNewOp<mlir::LLVM::ZeroOp>(zero, ty);
3553 return mlir::success();
3564 mlir::ConversionPatternRewriter &rewriter) const override {
3565 rewriter.replaceOpWithNewOp<mlir::LLVM::UnreachableOp>(unreach);
3566 return mlir::success();
3572 /// %0 = llvm.mlir.constant(0 : i64)
3581 mlir::ConversionPatternRewriter &rewriter) const override {
3582 mlir::Type idxTy = lowerTy().indexType();
3583 mlir::Location loc = isPresent.getLoc();
3586 if (mlir::isa<fir::BoxCharType>(isPresent.getVal().getType())) {
3588 mlir::cast<mlir::LLVM::LLVMStructType>(ptr.getType());
3591 ptr = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, ptr, 0);
3593 mlir::LLVM::ConstantOp c0 =
3595 auto addr = rewriter.create<mlir::LLVM::PtrToIntOp>(loc, idxTy, ptr);
3596 rewriter.replaceOpWithNewOp<mlir::LLVM::ICmpOp>(
3597 isPresent, mlir::LLVM::ICmpPredicate::ne, addr, c0);
3599 return mlir::success();
3604 /// `fir.absent !fir.ref<i64>` --> `llvm.mlir.zero : !llvm.ptr<i64>`
3610 mlir::ConversionPatternRewriter &rewriter) const override {
3611 mlir::Type ty = convertType(absent.getType());
3612 rewriter.replaceOpWithNewOp<mlir::LLVM::ZeroOp>(absent, ty);
3613 return mlir::success();
3622 static inline mlir::LLVM::FastmathFlagsAttr getLLVMFMFAttr(OPTY op) {
3623 return mlir::LLVM::FastmathFlagsAttr::get(
3625 mlir::arith::convertArithFastMathFlagsToLLVM(op.getFastmath()));
3630 static mlir::LLVM::InsertValueOp
3631 complexSum(OPTY sumop, mlir::ValueRange opnds,
3632 mlir::ConversionPatternRewriter &rewriter,
3634 mlir::LLVM::FastmathFlagsAttr fmf = getLLVMFMFAttr(sumop);
3635 mlir::Value a = opnds[0];
3636 mlir::Value b = opnds[1];
3638 mlir::Type eleTy = lowering.convertType(getComplexEleTy(sumop.getType()));
3639 mlir::Type ty = lowering.convertType(sumop.getType());
3640 auto x0 = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, a, 0);
3641 auto y0 = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, a, 1);
3642 auto x1 = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, b, 0);
3643 auto y1 = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, b, 1);
3646 auto r0 = rewriter.create<mlir::LLVM::UndefOp>(loc, ty);
3647 auto r1 = rewriter.create<mlir::LLVM::InsertValueOp>(loc, r0, rx, 0);
3648 return rewriter.create<mlir::LLVM::InsertValueOp>(loc, r1, ry, 1);
3658 mlir::ConversionPatternRewriter &rewriter) const override {
3661 auto r = complexSum<mlir::LLVM::FAddOp>(addc, adaptor.getOperands(),
3664 return mlir::success();
3673 mlir::ConversionPatternRewriter &rewriter) const override {
3676 auto r = complexSum<mlir::LLVM::FSubOp>(subc, adaptor.getOperands(),
3679 return mlir::success();
3689 mlir::ConversionPatternRewriter &rewriter) const override {
3693 mlir::LLVM::FastmathFlagsAttr fmf = getLLVMFMFAttr(mulc);
3694 mlir::Value a = adaptor.getOperands()[0];
3695 mlir::Value b = adaptor.getOperands()[1];
3697 mlir::Type eleTy = convertType(getComplexEleTy(mulc.getType()));
3698 mlir::Type ty = convertType(mulc.getType());
3699 auto x0 = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, a, 0);
3700 auto y0 = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, a, 1);
3701 auto x1 = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, b, 0);
3702 auto y1 = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, b, 1);
3703 auto xx = rewriter.create<mlir::LLVM::FMulOp>(loc, eleTy, x0, x1, fmf);
3704 auto yx = rewriter.create<mlir::LLVM::FMulOp>(loc, eleTy, y0, x1, fmf);
3705 auto xy = rewriter.create<mlir::LLVM::FMulOp>(loc, eleTy, x0, y1, fmf);
3706 auto ri = rewriter.create<mlir::LLVM::FAddOp>(loc, eleTy, xy, yx, fmf);
3707 auto yy = rewriter.create<mlir::LLVM::FMulOp>(loc, eleTy, y0, y1, fmf);
3708 auto rr = rewriter.create<mlir::LLVM::FSubOp>(loc, eleTy, xx, yy, fmf);
3709 auto ra = rewriter.create<mlir::LLVM::UndefOp>(loc, ty);
3710 auto r1 = rewriter.create<mlir::LLVM::InsertValueOp>(loc, ra, rr, 0);
3711 auto r0 = rewriter.create<mlir::LLVM::InsertValueOp>(loc, r1, ri, 1);
3713 return mlir::success();
3723 mlir::ConversionPatternRewriter &rewriter) const override {
3728 mlir::LLVM::FastmathFlagsAttr fmf = getLLVMFMFAttr(divc);
3729 mlir::Value a = adaptor.getOperands()[0];
3730 mlir::Value b = adaptor.getOperands()[1];
3732 mlir::Type eleTy = convertType(getComplexEleTy(divc.getType()));
3733 mlir::Type ty = convertType(divc.getType());
3734 auto x0 = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, a, 0);
3735 auto y0 = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, a, 1);
3736 auto x1 = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, b, 0);
3737 auto y1 = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, b, 1);
3738 auto xx = rewriter.create<mlir::LLVM::FMulOp>(loc, eleTy, x0, x1, fmf);
3739 auto x1x1 = rewriter.create<mlir::LLVM::FMulOp>(loc, eleTy, x1, x1, fmf);
3740 auto yx = rewriter.create<mlir::LLVM::FMulOp>(loc, eleTy, y0, x1, fmf);
3741 auto xy = rewriter.create<mlir::LLVM::FMulOp>(loc, eleTy, x0, y1, fmf);
3742 auto yy = rewriter.create<mlir::LLVM::FMulOp>(loc, eleTy, y0, y1, fmf);
3743 auto y1y1 = rewriter.create<mlir::LLVM::FMulOp>(loc, eleTy, y1, y1, fmf);
3744 auto d = rewriter.create<mlir::LLVM::FAddOp>(loc, eleTy, x1x1, y1y1, fmf);
3745 auto rrn = rewriter.create<mlir::LLVM::FAddOp>(loc, eleTy, xx, yy, fmf);
3746 auto rin = rewriter.create<mlir::LLVM::FSubOp>(loc, eleTy, yx, xy, fmf);
3747 auto rr = rewriter.create<mlir::LLVM::FDivOp>(loc, eleTy, rrn, d, fmf);
3748 auto ri = rewriter.create<mlir::LLVM::FDivOp>(loc, eleTy, rin, d, fmf);
3749 auto ra = rewriter.create<mlir::LLVM::UndefOp>(loc, ty);
3750 auto r1 = rewriter.create<mlir::LLVM::InsertValueOp>(loc, ra, rr, 0);
3751 auto r0 = rewriter.create<mlir::LLVM::InsertValueOp>(loc, r1, ri, 1);
3753 return mlir::success();
3763 mlir::ConversionPatternRewriter &rewriter) const override {
3768 mlir::Value o0 = adaptor.getOperands()[0];
3769 auto rp = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, o0, 0);
3770 auto ip = rewriter.create<mlir::LLVM::ExtractValueOp>(loc, o0, 1);
3771 auto nrp = rewriter.create<mlir::LLVM::FNegOp>(loc, eleTy, rp);
3772 auto nip = rewriter.create<mlir::LLVM::FNegOp>(loc, eleTy, ip);
3773 auto r = rewriter.create<mlir::LLVM::InsertValueOp>(loc, o0, nrp, 0);
3774 rewriter.replaceOpWithNewOp<mlir::LLVM::InsertValueOp>(neg, r, nip, 1);
3775 return mlir::success();
3784 mlir::ConversionPatternRewriter &rewriter) const override {
3786 mlir::Type pty = ::getLlvmPtrType(boxOffset.getContext());
3787 mlir::Type boxType = fir::unwrapRefType(boxOffset.getBoxRef().getType());
3788 mlir::Type llvmBoxTy =
3789 lowerTy().convertBoxTypeAsStruct(mlir::cast<fir::BaseBoxType>(boxType));
3793 rewriter.replaceOpWithNewOp<mlir::LLVM::GEPOp>(
3795 llvm::ArrayRef<mlir::LLVM::GEPArg>{0, fieldId});
3796 return mlir::success();
3813 mlir::ConversionPatternRewriter &rewriter) const final {
3817 return mlir::success();
3841 : public mlir::OpRewritePattern<mlir::LLVM::CallOp> {
3846 matchAndRewrite(mlir::LLVM::CallOp op,
3847 mlir::PatternRewriter &rewriter) const override {
3852 op.setCalleeAttr(mlir::SymbolRefAttr::get(op.getContext(), "_hypotf"));
3855 return mlir::success();
3860 : public mlir::OpRewritePattern<mlir::LLVM::LLVMFuncOp> {
3865 matchAndRewrite(mlir::LLVM::LLVMFuncOp op,
3866 mlir::PatternRewriter &rewriter) const override {
3871 return mlir::success();
3880 /// MLIR pass is used to lower residual Std dialect to LLVM IR dialect.
3886 mlir::ModuleOp getModule() { return getOperation(); }
3918 mlir::OpPassManager mathConvertionPM("builtin.module");
3926 mathConvertionPM.addPass(mlir::createConvertMathToROCDL());
3931 mlir::ConvertMathToFuncsOptions mathToFuncsOptions{};
3934 mlir::createConvertMathToFuncs(mathToFuncsOptions));
3935 mathConvertionPM.addPass(mlir::createConvertComplexToStandardPass());
3939 mathConvertionPM.addNestedPass<mlir::func::FuncOp>(
3940 mlir::createConvertMathToLLVMPass());
3941 if (mlir::failed(runPipeline(mathConvertionPM, mod)))
3944 std::optional<mlir::DataLayout> dl =
3947 mlir::emitError(mod.getLoc(),
3958 mlir::RewritePatternSet pattern(context);
3960 mlir::populateFuncToLLVMConversionPatterns(typeConverter, pattern);
3961 mlir::populateOpenMPToLLVMConversionPatterns(typeConverter, pattern);
3962 mlir::arith::populateArithToLLVMConversionPatterns(typeConverter, pattern);
3963 mlir::cf::populateControlFlowToLLVMConversionPatterns(typeConverter,
3965 mlir::cf::populateAssertToLLVMConversionPattern(typeConverter, pattern);
3969 mlir::populateMathToLibmConversionPatterns(pattern);
3970 mlir::populateComplexToLLVMConversionPatterns(typeConverter, pattern);
3971 mlir::populateVectorToLLVMConversionPatterns(typeConverter, pattern);
3977 mlir::ConversionTarget target{*context};
3978 target.addLegalDialect<mlir::LLVM::LLVMDialect>();
3983 mlir::configureOpenMPToLLVMConversionLegality(target, typeConverter);
3984 target.addLegalDialect<mlir::omp::OpenMPDialect>();
3985 target.addLegalDialect<mlir::acc::OpenACCDialect>();
3986 target.addLegalDialect<mlir::gpu::GPUDialect>();
3989 target.addLegalOp<mlir::ModuleOp>();
3996 target.addDynamicallyLegalOp<mlir::LLVM::CallOp>(
3997 [](mlir::LLVM::CallOp op) {
4003 target.addDynamicallyLegalOp<mlir::LLVM::LLVMFuncOp>(
4004 [](mlir::LLVM::LLVMFuncOp op) {
4010 if (mlir::failed(mlir::applyFullConversion(getModule(), target,
4018 mlir::OpPassManager comdatPM("builtin.module");
4019 comdatPM.addPass(mlir::LLVM::createLLVMAddComdats());
4020 if (mlir::failed(runPipeline(comdatPM, mod)))
4031 : public mlir::PassWrapper<LLVMIRLoweringPass,
4032 mlir::OperationPass<mlir::ModuleOp>> {
4038 mlir::ModuleOp getModule() { return getOperation(); }
4044 if (auto llvmModule = mlir::translateModuleToLLVMIR(
4050 mlir::emitError(mlir::UnknownLoc::get(ctx), "could not emit LLVM-IR\n");
4061 std::unique_ptr<mlir::Pass> fir::createFIRToLLVMPass() {
4065 std::unique_ptr<mlir::Pass>
4070 std::unique_ptr<mlir::Pass>
4077 const fir::LLVMTypeConverter &converter, mlir::RewritePatternSet &patterns,