diff --git a/flang/lib/Optimizer/Transforms/AbstractResult.cpp b/flang/lib/Optimizer/Transforms/AbstractResult.cpp index 59e2eeb76c715..391cfe3ceb9a2 100644 --- a/flang/lib/Optimizer/Transforms/AbstractResult.cpp +++ b/flang/lib/Optimizer/Transforms/AbstractResult.cpp @@ -137,9 +137,9 @@ class CallConversion : public mlir::OpRewritePattern { auto buffer = saveResult.getMemref(); mlir::Value arg = buffer; if (mustEmboxResult(result.getType(), shouldBoxResult)) - arg = rewriter.create( - loc, argType, buffer, saveResult.getShape(), /*slice*/ mlir::Value{}, - saveResult.getTypeparams()); + arg = fir::EmboxOp::create(rewriter, loc, argType, buffer, + saveResult.getShape(), /*slice*/ mlir::Value{}, + saveResult.getTypeparams()); llvm::SmallVector newResultTypes; bool isResultBuiltinCPtr = fir::isa_builtin_cptr_type(result.getType()); @@ -155,8 +155,8 @@ class CallConversion : public mlir::OpRewritePattern { if (!isResultBuiltinCPtr) newOperands.emplace_back(arg); newOperands.append(op.getOperands().begin(), op.getOperands().end()); - newOp = rewriter.create(loc, *op.getCallee(), - newResultTypes, newOperands); + newOp = fir::CallOp::create(rewriter, loc, *op.getCallee(), + newResultTypes, newOperands); } else { // Indirect calls. llvm::SmallVector newInputTypes; @@ -169,13 +169,13 @@ class CallConversion : public mlir::OpRewritePattern { llvm::SmallVector newOperands; newOperands.push_back( - rewriter.create(loc, newFuncTy, op.getOperand(0))); + fir::ConvertOp::create(rewriter, loc, newFuncTy, op.getOperand(0))); if (!isResultBuiltinCPtr) newOperands.push_back(arg); newOperands.append(op.getOperands().begin() + 1, op.getOperands().end()); - newOp = rewriter.create(loc, mlir::SymbolRefAttr{}, - newResultTypes, newOperands); + newOp = fir::CallOp::create(rewriter, loc, mlir::SymbolRefAttr{}, + newResultTypes, newOperands); } } @@ -191,8 +191,8 @@ class CallConversion : public mlir::OpRewritePattern { passArgPos = rewriter.getI32IntegerAttr(*op.getPassArgPos() + passArgShift); // TODO: propagate argument and result attributes (need to be shifted). - newOp = rewriter.create( - loc, newResultTypes, rewriter.getStringAttr(op.getMethod()), + newOp = fir::DispatchOp::create( + rewriter, loc, newResultTypes, rewriter.getStringAttr(op.getMethod()), op.getOperands()[0], newOperands, passArgPos, /*arg_attrs=*/nullptr, /*res_attrs=*/nullptr, op.getProcedureAttrsAttr()); @@ -280,7 +280,7 @@ processReturnLikeOp(OpTy ret, mlir::Value newArg, // register pass, this is possible for fir.box results, or fir.record // with no length parameters. Simply store the result in the result // storage. at the return point. - rewriter.create(loc, resultValue, newArg); + fir::StoreOp::create(rewriter, loc, resultValue, newArg); rewriter.replaceOpWithNewOp(ret); } // Delete result old local storage if unused. @@ -337,8 +337,8 @@ class AddrOfOpConversion : public mlir::OpRewritePattern { newFuncTy = getCPtrFunctionType(oldFuncTy); else newFuncTy = getNewFunctionType(oldFuncTy, shouldBoxResult); - auto newAddrOf = rewriter.create(addrOf.getLoc(), newFuncTy, - addrOf.getSymbol()); + auto newAddrOf = fir::AddrOfOp::create(rewriter, addrOf.getLoc(), newFuncTy, + addrOf.getSymbol()); // Rather than converting all op a function pointer might transit through // (e.g calls, stores, loads, converts...), cast new type to the abstract // type. A conversion will be added when calling indirect calls of abstract @@ -397,7 +397,7 @@ class AbstractResultOpt if (mustEmboxResult(resultType, shouldBoxResult)) { auto bufferType = fir::ReferenceType::get(resultType); rewriter.setInsertionPointToStart(&func.front()); - newArg = rewriter.create(loc, bufferType, newArg); + newArg = fir::BoxAddrOp::create(rewriter, loc, bufferType, newArg); } patterns.insert(context, newArg); target.addDynamicallyLegalOp( diff --git a/flang/lib/Optimizer/Transforms/AffineDemotion.cpp b/flang/lib/Optimizer/Transforms/AffineDemotion.cpp index d45f855c9078e..f1c66a5bbcf8c 100644 --- a/flang/lib/Optimizer/Transforms/AffineDemotion.cpp +++ b/flang/lib/Optimizer/Transforms/AffineDemotion.cpp @@ -60,9 +60,10 @@ class AffineLoadConversion if (!maybeExpandedMap) return failure(); - auto coorOp = rewriter.create( - op.getLoc(), fir::ReferenceType::get(op.getResult().getType()), - adaptor.getMemref(), *maybeExpandedMap); + auto coorOp = fir::CoordinateOp::create( + rewriter, op.getLoc(), + fir::ReferenceType::get(op.getResult().getType()), adaptor.getMemref(), + *maybeExpandedMap); rewriter.replaceOpWithNewOp(op, coorOp.getResult()); return success(); @@ -83,8 +84,9 @@ class AffineStoreConversion if (!maybeExpandedMap) return failure(); - auto coorOp = rewriter.create( - op.getLoc(), fir::ReferenceType::get(op.getValueToStore().getType()), + auto coorOp = fir::CoordinateOp::create( + rewriter, op.getLoc(), + fir::ReferenceType::get(op.getValueToStore().getType()), adaptor.getMemref(), *maybeExpandedMap); rewriter.replaceOpWithNewOp(op, adaptor.getValue(), coorOp.getResult()); diff --git a/flang/lib/Optimizer/Transforms/AffinePromotion.cpp b/flang/lib/Optimizer/Transforms/AffinePromotion.cpp index ef82e400bea14..b032767eef6f0 100644 --- a/flang/lib/Optimizer/Transforms/AffinePromotion.cpp +++ b/flang/lib/Optimizer/Transforms/AffinePromotion.cpp @@ -366,8 +366,9 @@ static mlir::Type coordinateArrayElement(fir::ArrayCoorOp op) { static void populateIndexArgs(fir::ArrayCoorOp acoOp, fir::ShapeOp shape, SmallVectorImpl &indexArgs, mlir::PatternRewriter &rewriter) { - auto one = rewriter.create( - acoOp.getLoc(), rewriter.getIndexType(), rewriter.getIndexAttr(1)); + auto one = mlir::arith::ConstantOp::create(rewriter, acoOp.getLoc(), + rewriter.getIndexType(), + rewriter.getIndexAttr(1)); auto extents = shape.getExtents(); for (auto i = extents.begin(); i < extents.end(); i++) { indexArgs.push_back(one); @@ -379,8 +380,9 @@ static void populateIndexArgs(fir::ArrayCoorOp acoOp, fir::ShapeOp shape, static void populateIndexArgs(fir::ArrayCoorOp acoOp, fir::ShapeShiftOp shape, SmallVectorImpl &indexArgs, mlir::PatternRewriter &rewriter) { - auto one = rewriter.create( - acoOp.getLoc(), rewriter.getIndexType(), rewriter.getIndexAttr(1)); + auto one = mlir::arith::ConstantOp::create(rewriter, acoOp.getLoc(), + rewriter.getIndexType(), + rewriter.getIndexAttr(1)); auto extents = shape.getPairs(); for (auto i = extents.begin(); i < extents.end();) { indexArgs.push_back(*i++); @@ -422,13 +424,13 @@ createAffineOps(mlir::Value arrayRef, mlir::PatternRewriter &rewriter) { populateIndexArgs(acoOp, indexArgs, rewriter); - auto affineApply = rewriter.create( - acoOp.getLoc(), affineMap, indexArgs); + auto affineApply = affine::AffineApplyOp::create(rewriter, acoOp.getLoc(), + affineMap, indexArgs); auto arrayElementType = coordinateArrayElement(acoOp); auto newType = mlir::MemRefType::get({mlir::ShapedType::kDynamic}, arrayElementType); - auto arrayConvert = rewriter.create(acoOp.getLoc(), newType, - acoOp.getMemref()); + auto arrayConvert = fir::ConvertOp::create(rewriter, acoOp.getLoc(), newType, + acoOp.getMemref()); return std::make_pair(affineApply, arrayConvert); } @@ -495,7 +497,7 @@ class AffineLoopConversion : public mlir::OpRewritePattern { affineFor.getRegionIterArgs()); if (!results.empty()) { rewriter.setInsertionPointToEnd(affineFor.getBody()); - rewriter.create(resultOp->getLoc(), results); + affine::AffineYieldOp::create(rewriter, resultOp->getLoc(), results); } rewriter.finalizeOpModification(affineFor.getOperation()); @@ -525,8 +527,8 @@ class AffineLoopConversion : public mlir::OpRewritePattern { std::pair positiveConstantStep(fir::DoLoopOp op, int64_t step, mlir::PatternRewriter &rewriter) const { - auto affineFor = rewriter.create( - op.getLoc(), ValueRange(op.getLowerBound()), + auto affineFor = affine::AffineForOp::create( + rewriter, op.getLoc(), ValueRange(op.getLowerBound()), mlir::AffineMap::get(0, 1, mlir::getAffineSymbolExpr(0, op.getContext())), ValueRange(op.getUpperBound()), @@ -543,24 +545,24 @@ class AffineLoopConversion : public mlir::OpRewritePattern { auto step = mlir::getAffineSymbolExpr(2, op.getContext()); mlir::AffineMap upperBoundMap = mlir::AffineMap::get( 0, 3, (upperBound - lowerBound + step).floorDiv(step)); - auto genericUpperBound = rewriter.create( - op.getLoc(), upperBoundMap, + auto genericUpperBound = affine::AffineApplyOp::create( + rewriter, op.getLoc(), upperBoundMap, ValueRange({op.getLowerBound(), op.getUpperBound(), op.getStep()})); auto actualIndexMap = mlir::AffineMap::get( 1, 2, (lowerBound + mlir::getAffineDimExpr(0, op.getContext())) * mlir::getAffineSymbolExpr(1, op.getContext())); - auto affineFor = rewriter.create( - op.getLoc(), ValueRange(), + auto affineFor = affine::AffineForOp::create( + rewriter, op.getLoc(), ValueRange(), AffineMap::getConstantMap(0, op.getContext()), genericUpperBound.getResult(), mlir::AffineMap::get(0, 1, 1 + mlir::getAffineSymbolExpr(0, op.getContext())), 1, op.getIterOperands()); rewriter.setInsertionPointToStart(affineFor.getBody()); - auto actualIndex = rewriter.create( - op.getLoc(), actualIndexMap, + auto actualIndex = affine::AffineApplyOp::create( + rewriter, op.getLoc(), actualIndexMap, ValueRange( {affineFor.getInductionVar(), op.getLowerBound(), op.getStep()})); return std::make_pair(affineFor, actualIndex.getResult()); @@ -588,8 +590,8 @@ class AffineIfConversion : public mlir::OpRewritePattern { << "AffineIfConversion: couldn't calculate affine condition\n";); return failure(); } - auto affineIf = rewriter.create( - op.getLoc(), affineCondition.getIntegerSet(), + auto affineIf = affine::AffineIfOp::create( + rewriter, op.getLoc(), affineCondition.getIntegerSet(), affineCondition.getAffineArgs(), !op.getElseRegion().empty()); rewriter.startOpModification(affineIf); affineIf.getThenBlock()->getOperations().splice( diff --git a/flang/lib/Optimizer/Transforms/ArrayValueCopy.cpp b/flang/lib/Optimizer/Transforms/ArrayValueCopy.cpp index 8544d17f62248..247ba953f3265 100644 --- a/flang/lib/Optimizer/Transforms/ArrayValueCopy.cpp +++ b/flang/lib/Optimizer/Transforms/ArrayValueCopy.cpp @@ -856,7 +856,7 @@ static bool getAdjustedExtents(mlir::Location loc, auto idxTy = rewriter.getIndexType(); if (isAssumedSize(result)) { // Use slice information to compute the extent of the column. - auto one = rewriter.create(loc, 1); + auto one = mlir::arith::ConstantIndexOp::create(rewriter, loc, 1); mlir::Value size = one; if (mlir::Value sliceArg = arrLoad.getSlice()) { if (auto sliceOp = @@ -896,14 +896,14 @@ static mlir::Value getOrReadExtentsAndShapeOp( mlir::cast(dyn_cast_ptrOrBoxEleTy(boxTy)).getDimension(); auto idxTy = rewriter.getIndexType(); for (decltype(rank) dim = 0; dim < rank; ++dim) { - auto dimVal = rewriter.create(loc, dim); - auto dimInfo = rewriter.create(loc, idxTy, idxTy, idxTy, - arrLoad.getMemref(), dimVal); + auto dimVal = mlir::arith::ConstantIndexOp::create(rewriter, loc, dim); + auto dimInfo = BoxDimsOp::create(rewriter, loc, idxTy, idxTy, idxTy, + arrLoad.getMemref(), dimVal); result.emplace_back(dimInfo.getResult(1)); } if (!arrLoad.getShape()) { auto shapeType = ShapeType::get(rewriter.getContext(), rank); - return rewriter.create(loc, shapeType, result); + return ShapeOp::create(rewriter, loc, shapeType, result); } auto shiftOp = arrLoad.getShape().getDefiningOp(); auto shapeShiftType = ShapeShiftType::get(rewriter.getContext(), rank); @@ -912,8 +912,8 @@ static mlir::Value getOrReadExtentsAndShapeOp( shapeShiftOperands.push_back(lb); shapeShiftOperands.push_back(extent); } - return rewriter.create(loc, shapeShiftType, - shapeShiftOperands); + return ShapeShiftOp::create(rewriter, loc, shapeShiftType, + shapeShiftOperands); } copyUsingSlice = getAdjustedExtents(loc, rewriter, arrLoad, result, arrLoad.getShape()); @@ -952,13 +952,13 @@ static mlir::Value genCoorOp(mlir::PatternRewriter &rewriter, auto module = load->getParentOfType(); FirOpBuilder builder(rewriter, module); auto typeparams = getTypeParamsIfRawData(loc, builder, load, alloc.getType()); - mlir::Value result = rewriter.create( - loc, eleTy, alloc, shape, slice, + mlir::Value result = ArrayCoorOp::create( + rewriter, loc, eleTy, alloc, shape, slice, llvm::ArrayRef{originated}.take_front(dimension), typeparams); if (dimension < originated.size()) - result = rewriter.create( - loc, resTy, result, + result = fir::CoordinateOp::create( + rewriter, loc, resTy, result, llvm::ArrayRef{originated}.drop_front(dimension)); return result; } @@ -971,13 +971,13 @@ static mlir::Value getCharacterLen(mlir::Location loc, FirOpBuilder &builder, // The loaded array is an emboxed value. Get the CHARACTER length from // the box value. auto eleSzInBytes = - builder.create(loc, charLenTy, load.getMemref()); + BoxEleSizeOp::create(builder, loc, charLenTy, load.getMemref()); auto kindSize = builder.getKindMap().getCharacterBitsize(charTy.getFKind()); auto kindByteSize = builder.createIntegerConstant(loc, charLenTy, kindSize / 8); - return builder.create(loc, eleSzInBytes, - kindByteSize); + return mlir::arith::DivSIOp::create(builder, loc, eleSzInBytes, + kindByteSize); } // The loaded array is a (set of) unboxed values. If the CHARACTER's // length is not a constant, it must be provided as a type parameter to @@ -1003,11 +1003,11 @@ void genArrayCopy(mlir::Location loc, mlir::PatternRewriter &rewriter, auto idxTy = rewriter.getIndexType(); // Build loop nest from column to row. for (auto sh : llvm::reverse(extents)) { - auto ubi = rewriter.create(loc, idxTy, sh); - auto zero = rewriter.create(loc, 0); - auto one = rewriter.create(loc, 1); - auto ub = rewriter.create(loc, idxTy, ubi, one); - auto loop = rewriter.create(loc, zero, ub, one); + auto ubi = ConvertOp::create(rewriter, loc, idxTy, sh); + auto zero = mlir::arith::ConstantIndexOp::create(rewriter, loc, 0); + auto one = mlir::arith::ConstantIndexOp::create(rewriter, loc, 1); + auto ub = mlir::arith::SubIOp::create(rewriter, loc, idxTy, ubi, one); + auto loop = DoLoopOp::create(rewriter, loc, zero, ub, one); rewriter.setInsertionPointToStart(loop.getBody()); indices.push_back(loop.getInductionVar()); } @@ -1015,13 +1015,13 @@ void genArrayCopy(mlir::Location loc, mlir::PatternRewriter &rewriter, std::reverse(indices.begin(), indices.end()); auto module = arrLoad->getParentOfType(); FirOpBuilder builder(rewriter, module); - auto fromAddr = rewriter.create( - loc, getEleTy(src.getType()), src, shapeOp, + auto fromAddr = ArrayCoorOp::create( + rewriter, loc, getEleTy(src.getType()), src, shapeOp, CopyIn && copyUsingSlice ? sliceOp : mlir::Value{}, factory::originateIndices(loc, rewriter, src.getType(), shapeOp, indices), getTypeParamsIfRawData(loc, builder, arrLoad, src.getType())); - auto toAddr = rewriter.create( - loc, getEleTy(dst.getType()), dst, shapeOp, + auto toAddr = ArrayCoorOp::create( + rewriter, loc, getEleTy(dst.getType()), dst, shapeOp, !CopyIn && copyUsingSlice ? sliceOp : mlir::Value{}, factory::originateIndices(loc, rewriter, dst.getType(), shapeOp, indices), getTypeParamsIfRawData(loc, builder, arrLoad, dst.getType())); @@ -1093,15 +1093,16 @@ allocateArrayTemp(mlir::Location loc, mlir::PatternRewriter &rewriter, findNonconstantExtents(baseType, extents); llvm::SmallVector typeParams = genArrayLoadTypeParameters(loc, rewriter, load); - mlir::Value allocmem = rewriter.create( - loc, dyn_cast_ptrOrBoxEleTy(baseType), typeParams, nonconstantExtents); + mlir::Value allocmem = + AllocMemOp::create(rewriter, loc, dyn_cast_ptrOrBoxEleTy(baseType), + typeParams, nonconstantExtents); mlir::Type eleType = fir::unwrapSequenceType(fir::unwrapPassByRefType(baseType)); if (fir::isRecordWithAllocatableMember(eleType)) { // The allocatable component descriptors need to be set to a clean // deallocated status before anything is done with them. - mlir::Value box = rewriter.create( - loc, fir::BoxType::get(allocmem.getType()), allocmem, shape, + mlir::Value box = fir::EmboxOp::create( + rewriter, loc, fir::BoxType::get(allocmem.getType()), allocmem, shape, /*slice=*/mlir::Value{}, typeParams); auto module = load->getParentOfType(); FirOpBuilder builder(rewriter, module); @@ -1111,12 +1112,12 @@ allocateArrayTemp(mlir::Location loc, mlir::PatternRewriter &rewriter, auto cleanup = [=](mlir::PatternRewriter &r) { FirOpBuilder builder(r, module); runtime::genDerivedTypeDestroy(builder, loc, box); - r.create(loc, allocmem); + FreeMemOp::create(r, loc, allocmem); }; return {allocmem, cleanup}; } auto cleanup = [=](mlir::PatternRewriter &r) { - r.create(loc, allocmem); + FreeMemOp::create(r, loc, allocmem); }; return {allocmem, cleanup}; } @@ -1257,7 +1258,7 @@ class ArrayUpdateConversion : public ArrayUpdateConversionBase { if (auto inEleTy = dyn_cast_ptrEleTy(input.getType())) { emitFatalError(loc, "array_update on references not supported"); } else { - rewriter.create(loc, input, coor); + fir::StoreOp::create(rewriter, loc, input, coor); } }; auto lhsEltRefType = toRefType(update.getMerge().getType()); @@ -1368,7 +1369,7 @@ class ArrayAmendConversion : public mlir::OpRewritePattern { auto *op = amend.getOperation(); rewriter.setInsertionPoint(op); auto loc = amend.getLoc(); - auto undef = rewriter.create(loc, amend.getType()); + auto undef = UndefOp::create(rewriter, loc, amend.getType()); rewriter.replaceOp(amend, undef.getResult()); return mlir::success(); } diff --git a/flang/lib/Optimizer/Transforms/AssumedRankOpConversion.cpp b/flang/lib/Optimizer/Transforms/AssumedRankOpConversion.cpp index 6af1cb988a4c1..4c7b228eefeb5 100644 --- a/flang/lib/Optimizer/Transforms/AssumedRankOpConversion.cpp +++ b/flang/lib/Optimizer/Transforms/AssumedRankOpConversion.cpp @@ -88,8 +88,8 @@ class ReboxAssumedRankConv (fir::isPolymorphicType(oldBoxType) || (newEleType != oldBoxType.unwrapInnerType())) && !fir::isPolymorphicType(newBoxType)) { - newDtype = builder.create( - loc, mlir::TypeAttr::get(newDerivedType)); + newDtype = fir::TypeDescOp::create(builder, loc, + mlir::TypeAttr::get(newDerivedType)); } else { newDtype = builder.createNullConstant(loc); } @@ -103,7 +103,7 @@ class ReboxAssumedRankConv rebox.getBox(), newDtype, newAttribute, lowerBoundModifier); - mlir::Value descValue = builder.create(loc, tempDesc); + mlir::Value descValue = fir::LoadOp::create(builder, loc, tempDesc); mlir::Value castDesc = builder.createConvert(loc, newBoxType, descValue); rewriter.replaceOp(rebox, castDesc); return mlir::success(); diff --git a/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp b/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp index 2dd6950b34897..baa8e591ee162 100644 --- a/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp +++ b/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp @@ -68,25 +68,26 @@ struct CUFAddConstructor // Symbol reference to CUFRegisterAllocator. builder.setInsertionPointToEnd(mod.getBody()); - auto registerFuncOp = builder.create( - loc, RTNAME_STRING(CUFRegisterAllocator), funcTy); + auto registerFuncOp = mlir::LLVM::LLVMFuncOp::create( + builder, loc, RTNAME_STRING(CUFRegisterAllocator), funcTy); registerFuncOp.setVisibility(mlir::SymbolTable::Visibility::Private); auto cufRegisterAllocatorRef = mlir::SymbolRefAttr::get( mod.getContext(), RTNAME_STRING(CUFRegisterAllocator)); builder.setInsertionPointToEnd(mod.getBody()); // Create the constructor function that call CUFRegisterAllocator. - auto func = builder.create(loc, cudaFortranCtorName, - funcTy); + auto func = mlir::LLVM::LLVMFuncOp::create(builder, loc, + cudaFortranCtorName, funcTy); func.setLinkage(mlir::LLVM::Linkage::Internal); builder.setInsertionPointToStart(func.addEntryBlock(builder)); - builder.create(loc, funcTy, cufRegisterAllocatorRef); + mlir::LLVM::CallOp::create(builder, loc, funcTy, cufRegisterAllocatorRef); auto gpuMod = symTab.lookup(cudaDeviceModuleName); if (gpuMod) { auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get(ctx); - auto registeredMod = builder.create( - loc, llvmPtrTy, mlir::SymbolRefAttr::get(ctx, gpuMod.getName())); + auto registeredMod = cuf::RegisterModuleOp::create( + builder, loc, llvmPtrTy, + mlir::SymbolRefAttr::get(ctx, gpuMod.getName())); fir::LLVMTypeConverter typeConverter(mod, /*applyTBAA=*/false, /*forceUnifiedTBAATree=*/false, *dl); @@ -96,7 +97,8 @@ struct CUFAddConstructor auto kernelName = mlir::SymbolRefAttr::get( builder.getStringAttr(cudaDeviceModuleName), {mlir::SymbolRefAttr::get(builder.getContext(), func.getName())}); - builder.create(loc, kernelName, registeredMod); + cuf::RegisterKernelOp::create(builder, loc, kernelName, + registeredMod); } } @@ -140,19 +142,19 @@ struct CUFAddConstructor auto sizeVal = builder.createIntegerConstant(loc, idxTy, *size); // Global variable address - mlir::Value addr = builder.create( - loc, globalOp.resultType(), globalOp.getSymbol()); + mlir::Value addr = fir::AddrOfOp::create( + builder, loc, globalOp.resultType(), globalOp.getSymbol()); llvm::SmallVector args{fir::runtime::createArguments( builder, loc, fTy, registeredMod, addr, gblName, sizeVal)}; - builder.create(loc, func, args); + fir::CallOp::create(builder, loc, func, args); } break; default: break; } } } - builder.create(loc, mlir::ValueRange{}); + mlir::LLVM::ReturnOp::create(builder, loc, mlir::ValueRange{}); // Create the llvm.global_ctor with the function. // TODO: We might want to have a utility that retrieve it if already @@ -165,8 +167,8 @@ struct CUFAddConstructor llvm::SmallVector data; priorities.push_back(0); data.push_back(mlir::LLVM::ZeroAttr::get(mod.getContext())); - builder.create( - mod.getLoc(), builder.getArrayAttr(funcs), + mlir::LLVM::GlobalCtorsOp::create( + builder, mod.getLoc(), builder.getArrayAttr(funcs), builder.getI32ArrayAttr(priorities), builder.getArrayAttr(data)); } }; diff --git a/flang/lib/Optimizer/Transforms/CUFComputeSharedMemoryOffsetsAndSize.cpp b/flang/lib/Optimizer/Transforms/CUFComputeSharedMemoryOffsetsAndSize.cpp index f6381ef8a8a21..5e910f7da6472 100644 --- a/flang/lib/Optimizer/Transforms/CUFComputeSharedMemoryOffsetsAndSize.cpp +++ b/flang/lib/Optimizer/Transforms/CUFComputeSharedMemoryOffsetsAndSize.cpp @@ -93,10 +93,11 @@ struct CUFComputeSharedMemoryOffsetsAndSize mlir::Value dynSize = builder.createIntegerConstant(loc, idxTy, tySize); for (auto extent : sharedOp.getShape()) - dynSize = builder.create(loc, dynSize, extent); + dynSize = + mlir::arith::MulIOp::create(builder, loc, dynSize, extent); if (crtDynOffset) - crtDynOffset = - builder.create(loc, crtDynOffset, dynSize); + crtDynOffset = mlir::arith::AddIOp::create(builder, loc, + crtDynOffset, dynSize); else crtDynOffset = dynSize; @@ -142,9 +143,9 @@ struct CUFComputeSharedMemoryOffsetsAndSize fir::GlobalOp::getDataAttrAttrName(globalOpName), cuf::DataAttributeAttr::get(gpuMod.getContext(), cuf::DataAttribute::Shared))); - auto sharedMem = builder.create( - funcOp.getLoc(), sharedMemGlobalName, false, false, sharedMemType, - init, linkage, attrs); + auto sharedMem = fir::GlobalOp::create( + builder, funcOp.getLoc(), sharedMemGlobalName, false, false, + sharedMemType, init, linkage, attrs); sharedMem.setAlignment(alignment); } } diff --git a/flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp b/flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp index fe69ffa8350af..a40ed95391c3a 100644 --- a/flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp +++ b/flang/lib/Optimizer/Transforms/CUFGPUToLLVMConversion.cpp @@ -46,27 +46,28 @@ static mlir::Value createKernelArgArray(mlir::Location loc, auto structTy = mlir::LLVM::LLVMStructType::getLiteral(ctx, structTypes); auto ptrTy = mlir::LLVM::LLVMPointerType::get(rewriter.getContext()); mlir::Type i32Ty = rewriter.getI32Type(); - auto zero = rewriter.create( - loc, i32Ty, rewriter.getIntegerAttr(i32Ty, 0)); - auto one = rewriter.create( - loc, i32Ty, rewriter.getIntegerAttr(i32Ty, 1)); + auto zero = mlir::LLVM::ConstantOp::create(rewriter, loc, i32Ty, + rewriter.getIntegerAttr(i32Ty, 0)); + auto one = mlir::LLVM::ConstantOp::create(rewriter, loc, i32Ty, + rewriter.getIntegerAttr(i32Ty, 1)); mlir::Value argStruct = - rewriter.create(loc, ptrTy, structTy, one); - auto size = rewriter.create( - loc, i32Ty, rewriter.getIntegerAttr(i32Ty, structTypes.size())); + mlir::LLVM::AllocaOp::create(rewriter, loc, ptrTy, structTy, one); + auto size = mlir::LLVM::ConstantOp::create( + rewriter, loc, i32Ty, rewriter.getIntegerAttr(i32Ty, structTypes.size())); mlir::Value argArray = - rewriter.create(loc, ptrTy, ptrTy, size); + mlir::LLVM::AllocaOp::create(rewriter, loc, ptrTy, ptrTy, size); for (auto [i, arg] : llvm::enumerate(operands)) { - auto indice = rewriter.create( - loc, i32Ty, rewriter.getIntegerAttr(i32Ty, i)); - mlir::Value structMember = rewriter.create( - loc, ptrTy, structTy, argStruct, - mlir::ArrayRef({zero, indice})); - rewriter.create(loc, arg, structMember); - mlir::Value arrayMember = rewriter.create( - loc, ptrTy, ptrTy, argArray, mlir::ArrayRef({indice})); - rewriter.create(loc, structMember, arrayMember); + auto indice = mlir::LLVM::ConstantOp::create( + rewriter, loc, i32Ty, rewriter.getIntegerAttr(i32Ty, i)); + mlir::Value structMember = + LLVM::GEPOp::create(rewriter, loc, ptrTy, structTy, argStruct, + mlir::ArrayRef({zero, indice})); + LLVM::StoreOp::create(rewriter, loc, arg, structMember); + mlir::Value arrayMember = + LLVM::GEPOp::create(rewriter, loc, ptrTy, ptrTy, argArray, + mlir::ArrayRef({indice})); + LLVM::StoreOp::create(rewriter, loc, structMember, arrayMember); } return argArray; } @@ -94,8 +95,8 @@ struct GPULaunchKernelConversion mlir::Value dynamicMemorySize = op.getDynamicSharedMemorySize(); mlir::Type i32Ty = rewriter.getI32Type(); if (!dynamicMemorySize) - dynamicMemorySize = rewriter.create( - loc, i32Ty, rewriter.getIntegerAttr(i32Ty, 0)); + dynamicMemorySize = mlir::LLVM::ConstantOp::create( + rewriter, loc, i32Ty, rewriter.getIntegerAttr(i32Ty, 0)); mlir::Value kernelArgs = createKernelArgArray(loc, adaptor.getKernelOperands(), rewriter); @@ -108,17 +109,17 @@ struct GPULaunchKernelConversion if (!funcOp) return mlir::failure(); kernelPtr = - rewriter.create(loc, ptrTy, funcOp.getName()); + LLVM::AddressOfOp::create(rewriter, loc, ptrTy, funcOp.getName()); } else { kernelPtr = - rewriter.create(loc, ptrTy, kernel.getName()); + LLVM::AddressOfOp::create(rewriter, loc, ptrTy, kernel.getName()); } auto llvmIntPtrType = mlir::IntegerType::get( ctx, this->getTypeConverter()->getPointerBitwidth(0)); auto voidTy = mlir::LLVM::LLVMVoidType::get(ctx); - mlir::Value nullPtr = rewriter.create(loc, ptrTy); + mlir::Value nullPtr = LLVM::ZeroOp::create(rewriter, loc, ptrTy); if (op.hasClusterSize()) { auto funcOp = mod.lookupSymbol( @@ -134,8 +135,8 @@ struct GPULaunchKernelConversion if (!funcOp) { mlir::OpBuilder::InsertionGuard insertGuard(rewriter); rewriter.setInsertionPointToStart(mod.getBody()); - auto launchKernelFuncOp = rewriter.create( - loc, RTNAME_STRING(CUFLaunchClusterKernel), funcTy); + auto launchKernelFuncOp = mlir::LLVM::LLVMFuncOp::create( + rewriter, loc, RTNAME_STRING(CUFLaunchClusterKernel), funcTy); launchKernelFuncOp.setVisibility( mlir::SymbolTable::Visibility::Private); } @@ -148,8 +149,8 @@ struct GPULaunchKernelConversion stream = adaptor.getAsyncDependencies().front(); } - rewriter.create( - loc, funcTy, cufLaunchClusterKernel, + mlir::LLVM::CallOp::create( + rewriter, loc, funcTy, cufLaunchClusterKernel, mlir::ValueRange{kernelPtr, adaptor.getClusterSizeX(), adaptor.getClusterSizeY(), adaptor.getClusterSizeZ(), adaptor.getGridSizeX(), adaptor.getGridSizeY(), @@ -178,7 +179,7 @@ struct GPULaunchKernelConversion mlir::OpBuilder::InsertionGuard insertGuard(rewriter); rewriter.setInsertionPointToStart(mod.getBody()); auto launchKernelFuncOp = - rewriter.create(loc, fctName, funcTy); + mlir::LLVM::LLVMFuncOp::create(rewriter, loc, fctName, funcTy); launchKernelFuncOp.setVisibility( mlir::SymbolTable::Visibility::Private); } @@ -191,8 +192,8 @@ struct GPULaunchKernelConversion stream = adaptor.getAsyncDependencies().front(); } - rewriter.create( - loc, funcTy, cufLaunchKernel, + mlir::LLVM::CallOp::create( + rewriter, loc, funcTy, cufLaunchKernel, mlir::ValueRange{kernelPtr, adaptor.getGridSizeX(), adaptor.getGridSizeY(), adaptor.getGridSizeZ(), adaptor.getBlockSizeX(), adaptor.getBlockSizeY(), @@ -222,11 +223,11 @@ static mlir::Value createAddressOfOp(mlir::ConversionPatternRewriter &rewriter, auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get( rewriter.getContext(), mlir::NVVM::NVVMMemorySpace::kSharedMemorySpace); if (auto g = gpuMod.lookupSymbol(sharedGlobalName)) - return rewriter.create(loc, llvmPtrTy, - g.getSymName()); + return mlir::LLVM::AddressOfOp::create(rewriter, loc, llvmPtrTy, + g.getSymName()); if (auto g = gpuMod.lookupSymbol(sharedGlobalName)) - return rewriter.create(loc, llvmPtrTy, - g.getSymName()); + return mlir::LLVM::AddressOfOp::create(rewriter, loc, llvmPtrTy, + g.getSymName()); return {}; } @@ -255,13 +256,13 @@ struct CUFSharedMemoryOpConversion if (!sharedGlobalAddr) mlir::emitError(loc, "Could not find the shared global operation\n"); - auto castPtr = rewriter.create( - loc, mlir::LLVM::LLVMPointerType::get(rewriter.getContext()), + auto castPtr = mlir::LLVM::AddrSpaceCastOp::create( + rewriter, loc, mlir::LLVM::LLVMPointerType::get(rewriter.getContext()), sharedGlobalAddr); mlir::Type baseType = castPtr->getResultTypes().front(); llvm::SmallVector gepArgs = {op.getOffset()}; - mlir::Value shmemPtr = rewriter.create( - loc, baseType, rewriter.getI8Type(), castPtr, gepArgs); + mlir::Value shmemPtr = mlir::LLVM::GEPOp::create( + rewriter, loc, baseType, rewriter.getI8Type(), castPtr, gepArgs); rewriter.replaceOp(op, {shmemPtr}); return mlir::success(); } diff --git a/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp b/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp index 750569c126642..cd7d33091f345 100644 --- a/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp +++ b/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp @@ -89,7 +89,7 @@ static mlir::Value createConvertOp(mlir::PatternRewriter &rewriter, mlir::Location loc, mlir::Type toTy, mlir::Value val) { if (val.getType() != toTy) - return rewriter.create(loc, toTy, val); + return fir::ConvertOp::create(rewriter, loc, toTy, val); return val; } @@ -118,7 +118,7 @@ static mlir::LogicalResult convertOpToCall(OpTy op, errmsg = op.getErrmsg(); } else { mlir::Type boxNoneTy = fir::BoxType::get(builder.getNoneType()); - errmsg = builder.create(loc, boxNoneTy).getResult(); + errmsg = fir::AbsentOp::create(builder, loc, boxNoneTy).getResult(); } llvm::SmallVector args; if constexpr (std::is_same_v) { @@ -148,7 +148,7 @@ static mlir::LogicalResult convertOpToCall(OpTy op, fir::runtime::createArguments(builder, loc, fTy, op.getBox(), hasStat, errmsg, sourceFile, sourceLine); } - auto callOp = builder.create(loc, func, args); + auto callOp = fir::CallOp::create(builder, loc, func, args); rewriter.replaceOp(op, callOp); return mlir::success(); } @@ -301,10 +301,11 @@ struct CUFAllocOpConversion : public mlir::OpRewritePattern { if (inDeviceContext(op.getOperation())) { // In device context just replace the cuf.alloc operation with a fir.alloc // the cuf.free will be removed. - auto allocaOp = rewriter.create( - loc, op.getInType(), op.getUniqName() ? *op.getUniqName() : "", - op.getBindcName() ? *op.getBindcName() : "", op.getTypeparams(), - op.getShape()); + auto allocaOp = + fir::AllocaOp::create(rewriter, loc, op.getInType(), + op.getUniqName() ? *op.getUniqName() : "", + op.getBindcName() ? *op.getBindcName() : "", + op.getTypeparams(), op.getShape()); allocaOp->setAttr(cuf::getDataAttrName(), op.getDataAttrAttr()); rewriter.replaceOp(op, allocaOp); return mlir::success(); @@ -338,14 +339,15 @@ struct CUFAllocOpConversion : public mlir::OpRewritePattern { assert(!op.getShape().empty() && "expect shape with dynamic arrays"); nbElem = builder.loadIfRef(loc, op.getShape()[0]); for (unsigned i = 1; i < op.getShape().size(); ++i) { - nbElem = rewriter.create( - loc, nbElem, builder.loadIfRef(loc, op.getShape()[i])); + nbElem = mlir::arith::MulIOp::create( + rewriter, loc, nbElem, + builder.loadIfRef(loc, op.getShape()[i])); } } else { nbElem = builder.createIntegerConstant(loc, builder.getIndexType(), seqTy.getConstantArraySize()); } - bytes = rewriter.create(loc, nbElem, width); + bytes = mlir::arith::MulIOp::create(rewriter, loc, nbElem, width); } else if (fir::isa_derived(op.getInType())) { mlir::Type structTy = typeConverter->convertType(op.getInType()); std::size_t structSize = dl->getTypeSizeInBits(structTy) / 8; @@ -363,7 +365,7 @@ struct CUFAllocOpConversion : public mlir::OpRewritePattern { loc, builder.getI32Type(), getMemType(op.getDataAttr())); llvm::SmallVector args{fir::runtime::createArguments( builder, loc, fTy, bytes, memTy, sourceFile, sourceLine)}; - auto callOp = builder.create(loc, func, args); + auto callOp = fir::CallOp::create(builder, loc, func, args); callOp->setAttr(cuf::getDataAttrName(), op.getDataAttrAttr()); auto convOp = builder.createConvert(loc, op.getResult().getType(), callOp.getResult(0)); @@ -386,7 +388,7 @@ struct CUFAllocOpConversion : public mlir::OpRewritePattern { llvm::SmallVector args{fir::runtime::createArguments( builder, loc, fTy, sizeInBytes, sourceFile, sourceLine)}; - auto callOp = builder.create(loc, func, args); + auto callOp = fir::CallOp::create(builder, loc, func, args); callOp->setAttr(cuf::getDataAttrName(), op.getDataAttrAttr()); auto convOp = builder.createConvert(loc, op.getResult().getType(), callOp.getResult(0)); @@ -414,8 +416,9 @@ struct CUFDeviceAddressOpConversion op.getHostSymbol().getRootReference().getValue())) { auto mod = op->getParentOfType(); mlir::Location loc = op.getLoc(); - auto hostAddr = rewriter.create( - loc, fir::ReferenceType::get(global.getType()), op.getHostSymbol()); + auto hostAddr = fir::AddrOfOp::create( + rewriter, loc, fir::ReferenceType::get(global.getType()), + op.getHostSymbol()); fir::FirOpBuilder builder(rewriter, mod); mlir::func::FuncOp callee = fir::runtime::getRuntimeFunc(loc, @@ -428,7 +431,7 @@ struct CUFDeviceAddressOpConversion fir::factory::locationToLineNo(builder, loc, fTy.getInput(2)); llvm::SmallVector args{fir::runtime::createArguments( builder, loc, fTy, conv, sourceFile, sourceLine)}; - auto call = rewriter.create(loc, callee, args); + auto call = fir::CallOp::create(rewriter, loc, callee, args); mlir::Value addr = createConvertOp(rewriter, loc, hostAddr.getType(), call->getResult(0)); rewriter.replaceOp(op, addr.getDefiningOp()); @@ -456,8 +459,8 @@ struct DeclareOpConversion : public mlir::OpRewritePattern { addrOfOp.getSymbol().getRootReference().getValue())) { if (cuf::isRegisteredDeviceGlobal(global)) { rewriter.setInsertionPointAfter(addrOfOp); - mlir::Value devAddr = rewriter.create( - op.getLoc(), addrOfOp.getType(), addrOfOp.getSymbol()); + mlir::Value devAddr = cuf::DeviceAddressOp::create( + rewriter, op.getLoc(), addrOfOp.getType(), addrOfOp.getSymbol()); rewriter.startOpModification(op); op.getMemrefMutable().assign(devAddr); rewriter.finalizeOpModification(op); @@ -502,7 +505,7 @@ struct CUFFreeOpConversion : public mlir::OpRewritePattern { loc, builder.getI32Type(), getMemType(op.getDataAttr())); llvm::SmallVector args{fir::runtime::createArguments( builder, loc, fTy, op.getDevptr(), memTy, sourceFile, sourceLine)}; - builder.create(loc, func, args); + fir::CallOp::create(builder, loc, func, args); rewriter.eraseOp(op); return mlir::success(); } @@ -515,7 +518,7 @@ struct CUFFreeOpConversion : public mlir::OpRewritePattern { fir::factory::locationToLineNo(builder, loc, fTy.getInput(2)); llvm::SmallVector args{fir::runtime::createArguments( builder, loc, fTy, op.getDevptr(), sourceFile, sourceLine)}; - auto callOp = builder.create(loc, func, args); + auto callOp = fir::CallOp::create(builder, loc, func, args); callOp->setAttr(cuf::getDataAttrName(), op.getDataAttrAttr()); rewriter.eraseOp(op); return mlir::success(); @@ -558,18 +561,18 @@ static mlir::Value emboxSrc(mlir::PatternRewriter &rewriter, srcTy = fir::LogicalType::get(rewriter.getContext(), 4); src = createConvertOp(rewriter, loc, srcTy, src); addr = builder.createTemporary(loc, srcTy); - builder.create(loc, src, addr); + fir::StoreOp::create(builder, loc, src, addr); } else { if (dstEleTy && fir::isa_trivial(dstEleTy) && srcTy != dstEleTy) { // Use dstEleTy and convert to avoid assign mismatch. addr = builder.createTemporary(loc, dstEleTy); - auto conv = builder.create(loc, dstEleTy, src); - builder.create(loc, conv, addr); + auto conv = fir::ConvertOp::create(builder, loc, dstEleTy, src); + fir::StoreOp::create(builder, loc, conv, addr); srcTy = dstEleTy; } else { // Put constant in memory if it is not. addr = builder.createTemporary(loc, srcTy); - builder.create(loc, src, addr); + fir::StoreOp::create(builder, loc, src, addr); } } } else { @@ -582,7 +585,7 @@ static mlir::Value emboxSrc(mlir::PatternRewriter &rewriter, /*slice=*/nullptr, lenParams, /*tdesc=*/nullptr); mlir::Value src = builder.createTemporary(loc, box.getType()); - builder.create(loc, box, src); + fir::StoreOp::create(builder, loc, box, src); return src; } @@ -601,7 +604,7 @@ static mlir::Value emboxDst(mlir::PatternRewriter &rewriter, /*slice=*/nullptr, lenParams, /*tdesc=*/nullptr); mlir::Value dst = builder.createTemporary(loc, dstBox.getType()); - builder.create(loc, dstBox, dst); + fir::StoreOp::create(builder, loc, dstBox, dst); return dst; } @@ -660,7 +663,7 @@ struct CUFDataTransferOpConversion fir::factory::locationToLineNo(builder, loc, fTy.getInput(4)); llvm::SmallVector args{fir::runtime::createArguments( builder, loc, fTy, dst, src, modeValue, sourceFile, sourceLine)}; - builder.create(loc, func, args); + fir::CallOp::create(builder, loc, func, args); rewriter.eraseOp(op); return mlir::success(); } @@ -679,12 +682,12 @@ struct CUFDataTransferOpConversion extents.push_back(i.value()); } - nbElement = rewriter.create(loc, i64Ty, extents[0]); + nbElement = fir::ConvertOp::create(rewriter, loc, i64Ty, extents[0]); for (unsigned i = 1; i < extents.size(); ++i) { auto operand = - rewriter.create(loc, i64Ty, extents[i]); + fir::ConvertOp::create(rewriter, loc, i64Ty, extents[i]); nbElement = - rewriter.create(loc, nbElement, operand); + mlir::arith::MulIOp::create(rewriter, loc, nbElement, operand); } } else { if (auto seqTy = mlir::dyn_cast_or_null(dstTy)) @@ -699,12 +702,11 @@ struct CUFDataTransferOpConversion } else { width = computeWidth(loc, dstTy, kindMap); } - mlir::Value widthValue = rewriter.create( - loc, i64Ty, rewriter.getIntegerAttr(i64Ty, width)); - mlir::Value bytes = - nbElement - ? rewriter.create(loc, nbElement, widthValue) - : widthValue; + mlir::Value widthValue = mlir::arith::ConstantOp::create( + rewriter, loc, i64Ty, rewriter.getIntegerAttr(i64Ty, width)); + mlir::Value bytes = nbElement ? mlir::arith::MulIOp::create( + rewriter, loc, nbElement, widthValue) + : widthValue; mlir::func::FuncOp func = fir::runtime::getRuntimeFunc(loc, @@ -719,13 +721,13 @@ struct CUFDataTransferOpConversion // Materialize the src if constant. if (matchPattern(src.getDefiningOp(), mlir::m_Constant())) { mlir::Value temp = builder.createTemporary(loc, srcTy); - builder.create(loc, src, temp); + fir::StoreOp::create(builder, loc, src, temp); src = temp; } llvm::SmallVector args{ fir::runtime::createArguments(builder, loc, fTy, dst, src, bytes, modeValue, sourceFile, sourceLine)}; - builder.create(loc, func, args); + fir::CallOp::create(builder, loc, func, args); rewriter.eraseOp(op); return mlir::success(); } @@ -734,7 +736,7 @@ struct CUFDataTransferOpConversion if (mlir::isa(val.getDefiningOp())) { // Materialize the box to memory to be able to call the runtime. mlir::Value box = builder.createTemporary(loc, val.getType()); - builder.create(loc, val, box); + fir::StoreOp::create(builder, loc, val, box); return box; } return val; @@ -768,7 +770,7 @@ struct CUFDataTransferOpConversion fir::factory::locationToLineNo(builder, loc, fTy.getInput(4)); llvm::SmallVector args{fir::runtime::createArguments( builder, loc, fTy, dst, src, modeValue, sourceFile, sourceLine)}; - builder.create(loc, func, args); + fir::CallOp::create(builder, loc, func, args); rewriter.eraseOp(op); } else { // Transfer from a descriptor. @@ -784,7 +786,7 @@ struct CUFDataTransferOpConversion fir::factory::locationToLineNo(builder, loc, fTy.getInput(4)); llvm::SmallVector args{fir::runtime::createArguments( builder, loc, fTy, dst, src, modeValue, sourceFile, sourceLine)}; - builder.create(loc, func, args); + fir::CallOp::create(builder, loc, func, args); rewriter.eraseOp(op); } return mlir::success(); @@ -810,20 +812,21 @@ struct CUFLaunchOpConversion mlir::PatternRewriter &rewriter) const override { mlir::Location loc = op.getLoc(); auto idxTy = mlir::IndexType::get(op.getContext()); - mlir::Value zero = rewriter.create( - loc, rewriter.getIntegerType(32), rewriter.getI32IntegerAttr(0)); + mlir::Value zero = mlir::arith::ConstantOp::create( + rewriter, loc, rewriter.getIntegerType(32), + rewriter.getI32IntegerAttr(0)); auto gridSizeX = - rewriter.create(loc, idxTy, op.getGridX()); + mlir::arith::IndexCastOp::create(rewriter, loc, idxTy, op.getGridX()); auto gridSizeY = - rewriter.create(loc, idxTy, op.getGridY()); + mlir::arith::IndexCastOp::create(rewriter, loc, idxTy, op.getGridY()); auto gridSizeZ = - rewriter.create(loc, idxTy, op.getGridZ()); + mlir::arith::IndexCastOp::create(rewriter, loc, idxTy, op.getGridZ()); auto blockSizeX = - rewriter.create(loc, idxTy, op.getBlockX()); + mlir::arith::IndexCastOp::create(rewriter, loc, idxTy, op.getBlockX()); auto blockSizeY = - rewriter.create(loc, idxTy, op.getBlockY()); + mlir::arith::IndexCastOp::create(rewriter, loc, idxTy, op.getBlockY()); auto blockSizeZ = - rewriter.create(loc, idxTy, op.getBlockZ()); + mlir::arith::IndexCastOp::create(rewriter, loc, idxTy, op.getBlockZ()); auto kernelName = mlir::SymbolRefAttr::get( rewriter.getStringAttr(cudaDeviceModuleName), {mlir::SymbolRefAttr::get( @@ -835,12 +838,12 @@ struct CUFLaunchOpConversion op.getCallee().getLeafReference())) { if (auto clusterDimsAttr = funcOp->getAttrOfType( cuf::getClusterDimsAttrName())) { - clusterDimX = rewriter.create( - loc, clusterDimsAttr.getX().getInt()); - clusterDimY = rewriter.create( - loc, clusterDimsAttr.getY().getInt()); - clusterDimZ = rewriter.create( - loc, clusterDimsAttr.getZ().getInt()); + clusterDimX = mlir::arith::ConstantIndexOp::create( + rewriter, loc, clusterDimsAttr.getX().getInt()); + clusterDimY = mlir::arith::ConstantIndexOp::create( + rewriter, loc, clusterDimsAttr.getY().getInt()); + clusterDimZ = mlir::arith::ConstantIndexOp::create( + rewriter, loc, clusterDimsAttr.getZ().getInt()); } procAttr = funcOp->getAttrOfType(cuf::getProcAttrName()); @@ -870,8 +873,9 @@ struct CUFLaunchOpConversion args.push_back(arg); } mlir::Value dynamicShmemSize = op.getBytes() ? op.getBytes() : zero; - auto gpuLaunchOp = rewriter.create( - loc, kernelName, mlir::gpu::KernelDim3{gridSizeX, gridSizeY, gridSizeZ}, + auto gpuLaunchOp = mlir::gpu::LaunchFuncOp::create( + rewriter, loc, kernelName, + mlir::gpu::KernelDim3{gridSizeX, gridSizeY, gridSizeZ}, mlir::gpu::KernelDim3{blockSizeX, blockSizeY, blockSizeZ}, dynamicShmemSize, args); if (clusterDimX && clusterDimY && clusterDimZ) { @@ -883,7 +887,7 @@ struct CUFLaunchOpConversion mlir::OpBuilder::InsertionGuard guard(rewriter); rewriter.setInsertionPoint(gpuLaunchOp); mlir::Value stream = - rewriter.create(loc, op.getStream()); + cuf::StreamCastOp::create(rewriter, loc, op.getStream()); gpuLaunchOp.getAsyncDependenciesMutable().append(stream); } if (procAttr) @@ -916,8 +920,9 @@ struct CUFSyncDescriptorOpConversion if (!globalOp) return mlir::failure(); - auto hostAddr = builder.create( - loc, fir::ReferenceType::get(globalOp.getType()), op.getGlobalName()); + auto hostAddr = fir::AddrOfOp::create( + builder, loc, fir::ReferenceType::get(globalOp.getType()), + op.getGlobalName()); fir::runtime::cuda::genSyncGlobalDescriptor(builder, loc, hostAddr); op.erase(); return mlir::success(); diff --git a/flang/lib/Optimizer/Transforms/CharacterConversion.cpp b/flang/lib/Optimizer/Transforms/CharacterConversion.cpp index aee7e8ca5cb66..13da38e92c234 100644 --- a/flang/lib/Optimizer/Transforms/CharacterConversion.cpp +++ b/flang/lib/Optimizer/Transforms/CharacterConversion.cpp @@ -48,12 +48,13 @@ class CharacterConvertConversion << "running character conversion on " << conv << '\n'); // Establish a loop that executes count iterations. - auto zero = rewriter.create(loc, 0); - auto one = rewriter.create(loc, 1); + auto zero = mlir::arith::ConstantIndexOp::create(rewriter, loc, 0); + auto one = mlir::arith::ConstantIndexOp::create(rewriter, loc, 1); auto idxTy = rewriter.getIndexType(); - auto castCnt = rewriter.create(loc, idxTy, conv.getCount()); - auto countm1 = rewriter.create(loc, castCnt, one); - auto loop = rewriter.create(loc, zero, countm1, one); + auto castCnt = + fir::ConvertOp::create(rewriter, loc, idxTy, conv.getCount()); + auto countm1 = mlir::arith::SubIOp::create(rewriter, loc, castCnt, one); + auto loop = fir::DoLoopOp::create(rewriter, loc, zero, countm1, one); auto insPt = rewriter.saveInsertionPoint(); rewriter.setInsertionPointToStart(loop.getBody()); @@ -75,21 +76,22 @@ class CharacterConvertConversion auto toTy = rewriter.getIntegerType(toBits); auto toPtrTy = pointerType(toBits); auto fromPtr = - rewriter.create(loc, fromPtrTy, conv.getFrom()); - auto toPtr = rewriter.create(loc, toPtrTy, conv.getTo()); + fir::ConvertOp::create(rewriter, loc, fromPtrTy, conv.getFrom()); + auto toPtr = fir::ConvertOp::create(rewriter, loc, toPtrTy, conv.getTo()); auto getEleTy = [&](unsigned bits) { return fir::ReferenceType::get(rewriter.getIntegerType(bits)); }; - auto fromi = rewriter.create( - loc, getEleTy(fromBits), fromPtr, - mlir::ValueRange{loop.getInductionVar()}); - auto toi = rewriter.create( - loc, getEleTy(toBits), toPtr, mlir::ValueRange{loop.getInductionVar()}); - auto load = rewriter.create(loc, fromi); + auto fromi = + fir::CoordinateOp::create(rewriter, loc, getEleTy(fromBits), fromPtr, + mlir::ValueRange{loop.getInductionVar()}); + auto toi = + fir::CoordinateOp::create(rewriter, loc, getEleTy(toBits), toPtr, + mlir::ValueRange{loop.getInductionVar()}); + auto load = fir::LoadOp::create(rewriter, loc, fromi); mlir::Value icast = (fromBits >= toBits) - ? rewriter.create(loc, toTy, load).getResult() - : rewriter.create(loc, toTy, load) + ? fir::ConvertOp::create(rewriter, loc, toTy, load).getResult() + : mlir::arith::ExtUIOp::create(rewriter, loc, toTy, load) .getResult(); rewriter.replaceOpWithNewOp(conv, icast, toi); rewriter.restoreInsertionPoint(insPt); diff --git a/flang/lib/Optimizer/Transforms/ConstantArgumentGlobalisation.cpp b/flang/lib/Optimizer/Transforms/ConstantArgumentGlobalisation.cpp index 239a7cdaa4cf2..afafbd8179aff 100644 --- a/flang/lib/Optimizer/Transforms/ConstantArgumentGlobalisation.cpp +++ b/flang/lib/Optimizer/Transforms/ConstantArgumentGlobalisation.cpp @@ -111,11 +111,11 @@ class CallOpRewriter : public mlir::OpRewritePattern { builder.insert(cln); mlir::Value val = builder.createConvert(loc, varTy, cln->getResult(0)); - builder.create(loc, val); + fir::HasValueOp::create(builder, loc, val); }, builder.createInternalLinkage()); - mlir::Value addr = builder.create(loc, global.resultType(), - global.getSymbol()); + mlir::Value addr = fir::AddrOfOp::create( + builder, loc, global.resultType(), global.getSymbol()); newOperands.push_back(addr); needUpdate = true; } @@ -125,11 +125,11 @@ class CallOpRewriter : public mlir::OpRewritePattern { llvm::SmallVector newResultTypes; newResultTypes.append(callOp.getResultTypes().begin(), callOp.getResultTypes().end()); - fir::CallOp newOp = builder.create( - loc, - callOp.getCallee().has_value() ? callOp.getCallee().value() - : mlir::SymbolRefAttr{}, - newResultTypes, newOperands); + fir::CallOp newOp = fir::CallOp::create(builder, loc, + callOp.getCallee().has_value() + ? callOp.getCallee().value() + : mlir::SymbolRefAttr{}, + newResultTypes, newOperands); // Copy all the attributes from the old to new op. newOp->setAttrs(callOp->getAttrs()); rewriter.replaceOp(callOp, newOp); diff --git a/flang/lib/Optimizer/Transforms/ControlFlowConverter.cpp b/flang/lib/Optimizer/Transforms/ControlFlowConverter.cpp index 3d35803e6a2d3..e466aed753e63 100644 --- a/flang/lib/Optimizer/Transforms/ControlFlowConverter.cpp +++ b/flang/lib/Optimizer/Transforms/ControlFlowConverter.cpp @@ -83,17 +83,17 @@ class CfgLoopConv : public mlir::OpRewritePattern { // Initalization block rewriter.setInsertionPointToEnd(initBlock); - auto diff = rewriter.create(loc, high, low); - auto distance = rewriter.create(loc, diff, step); + auto diff = mlir::arith::SubIOp::create(rewriter, loc, high, low); + auto distance = mlir::arith::AddIOp::create(rewriter, loc, diff, step); mlir::Value iters = - rewriter.create(loc, distance, step); + mlir::arith::DivSIOp::create(rewriter, loc, distance, step); if (forceLoopToExecuteOnce) { - auto zero = rewriter.create(loc, 0); - auto cond = rewriter.create( - loc, arith::CmpIPredicate::sle, iters, zero); - auto one = rewriter.create(loc, 1); - iters = rewriter.create(loc, cond, one, iters); + auto zero = mlir::arith::ConstantIndexOp::create(rewriter, loc, 0); + auto cond = mlir::arith::CmpIOp::create( + rewriter, loc, arith::CmpIPredicate::sle, iters, zero); + auto one = mlir::arith::ConstantIndexOp::create(rewriter, loc, 1); + iters = mlir::arith::SelectOp::create(rewriter, loc, cond, one, iters); } llvm::SmallVector loopOperands; @@ -102,20 +102,20 @@ class CfgLoopConv : public mlir::OpRewritePattern { loopOperands.append(operands.begin(), operands.end()); loopOperands.push_back(iters); - rewriter.create(loc, conditionalBlock, loopOperands); + mlir::cf::BranchOp::create(rewriter, loc, conditionalBlock, loopOperands); // Last loop block auto *terminator = lastBlock->getTerminator(); rewriter.setInsertionPointToEnd(lastBlock); auto iv = conditionalBlock->getArgument(0); mlir::Value steppedIndex = - rewriter.create(loc, iv, step, iofAttr); + mlir::arith::AddIOp::create(rewriter, loc, iv, step, iofAttr); assert(steppedIndex && "must be a Value"); auto lastArg = conditionalBlock->getNumArguments() - 1; auto itersLeft = conditionalBlock->getArgument(lastArg); - auto one = rewriter.create(loc, 1); + auto one = mlir::arith::ConstantIndexOp::create(rewriter, loc, 1); mlir::Value itersMinusOne = - rewriter.create(loc, itersLeft, one); + mlir::arith::SubIOp::create(rewriter, loc, itersLeft, one); llvm::SmallVector loopCarried; loopCarried.push_back(steppedIndex); @@ -123,8 +123,8 @@ class CfgLoopConv : public mlir::OpRewritePattern { : terminator->operand_begin(); loopCarried.append(begin, terminator->operand_end()); loopCarried.push_back(itersMinusOne); - auto backEdge = - rewriter.create(loc, conditionalBlock, loopCarried); + auto backEdge = mlir::cf::BranchOp::create(rewriter, loc, conditionalBlock, + loopCarried); rewriter.eraseOp(terminator); // Copy loop annotations from the do loop to the loop back edge. @@ -133,13 +133,13 @@ class CfgLoopConv : public mlir::OpRewritePattern { // Conditional block rewriter.setInsertionPointToEnd(conditionalBlock); - auto zero = rewriter.create(loc, 0); - auto comparison = rewriter.create( - loc, arith::CmpIPredicate::sgt, itersLeft, zero); + auto zero = mlir::arith::ConstantIndexOp::create(rewriter, loc, 0); + auto comparison = mlir::arith::CmpIOp::create( + rewriter, loc, arith::CmpIPredicate::sgt, itersLeft, zero); - rewriter.create( - loc, comparison, firstBlock, llvm::ArrayRef(), endBlock, - llvm::ArrayRef()); + mlir::cf::CondBranchOp::create(rewriter, loc, comparison, firstBlock, + llvm::ArrayRef(), endBlock, + llvm::ArrayRef()); // The result of the loop operation is the values of the condition block // arguments except the induction variable on the last iteration. @@ -180,7 +180,7 @@ class CfgIfConv : public mlir::OpRewritePattern { continueBlock = rewriter.createBlock( remainingOpsBlock, ifOp.getResultTypes(), llvm::SmallVector(ifOp.getNumResults(), loc)); - rewriter.create(loc, remainingOpsBlock); + mlir::cf::BranchOp::create(rewriter, loc, remainingOpsBlock); } // Move blocks from the "then" region to the region containing 'fir.if', @@ -190,8 +190,8 @@ class CfgIfConv : public mlir::OpRewritePattern { auto *ifOpTerminator = ifOpRegion.back().getTerminator(); auto ifOpTerminatorOperands = ifOpTerminator->getOperands(); rewriter.setInsertionPointToEnd(&ifOpRegion.back()); - rewriter.create(loc, continueBlock, - ifOpTerminatorOperands); + mlir::cf::BranchOp::create(rewriter, loc, continueBlock, + ifOpTerminatorOperands); rewriter.eraseOp(ifOpTerminator); rewriter.inlineRegionBefore(ifOpRegion, continueBlock); @@ -205,16 +205,17 @@ class CfgIfConv : public mlir::OpRewritePattern { auto *otherwiseTerm = otherwiseRegion.back().getTerminator(); auto otherwiseTermOperands = otherwiseTerm->getOperands(); rewriter.setInsertionPointToEnd(&otherwiseRegion.back()); - rewriter.create(loc, continueBlock, - otherwiseTermOperands); + mlir::cf::BranchOp::create(rewriter, loc, continueBlock, + otherwiseTermOperands); rewriter.eraseOp(otherwiseTerm); rewriter.inlineRegionBefore(otherwiseRegion, continueBlock); } rewriter.setInsertionPointToEnd(condBlock); - auto branchOp = rewriter.create( - loc, ifOp.getCondition(), ifOpBlock, llvm::ArrayRef(), - otherwiseBlock, llvm::ArrayRef()); + auto branchOp = mlir::cf::CondBranchOp::create( + rewriter, loc, ifOp.getCondition(), ifOpBlock, + llvm::ArrayRef(), otherwiseBlock, + llvm::ArrayRef()); llvm::ArrayRef weights = ifOp.getWeights(); if (!weights.empty()) branchOp.setWeights(weights); @@ -269,7 +270,7 @@ class CfgIterWhileConv : public mlir::OpRewritePattern { rewriter.setInsertionPointToEnd(lastBodyBlock); auto step = whileOp.getStep(); mlir::Value stepped = - rewriter.create(loc, iv, step, iofAttr); + mlir::arith::AddIOp::create(rewriter, loc, iv, step, iofAttr); assert(stepped && "must be a Value"); llvm::SmallVector loopCarried; @@ -278,7 +279,7 @@ class CfgIterWhileConv : public mlir::OpRewritePattern { ? std::next(terminator->operand_begin()) : terminator->operand_begin(); loopCarried.append(begin, terminator->operand_end()); - rewriter.create(loc, conditionBlock, loopCarried); + mlir::cf::BranchOp::create(rewriter, loc, conditionBlock, loopCarried); rewriter.eraseOp(terminator); // Compute loop bounds before branching to the condition. @@ -293,31 +294,31 @@ class CfgIterWhileConv : public mlir::OpRewritePattern { destOperands.push_back(lowerBound); auto iterOperands = whileOp.getIterOperands(); destOperands.append(iterOperands.begin(), iterOperands.end()); - rewriter.create(loc, conditionBlock, destOperands); + mlir::cf::BranchOp::create(rewriter, loc, conditionBlock, destOperands); // With the body block done, we can fill in the condition block. rewriter.setInsertionPointToEnd(conditionBlock); // The comparison depends on the sign of the step value. We fully expect // this expression to be folded by the optimizer or LLVM. This expression // is written this way so that `step == 0` always returns `false`. - auto zero = rewriter.create(loc, 0); - auto compl0 = rewriter.create( - loc, arith::CmpIPredicate::slt, zero, step); - auto compl1 = rewriter.create( - loc, arith::CmpIPredicate::sle, iv, upperBound); - auto compl2 = rewriter.create( - loc, arith::CmpIPredicate::slt, step, zero); - auto compl3 = rewriter.create( - loc, arith::CmpIPredicate::sle, upperBound, iv); - auto cmp0 = rewriter.create(loc, compl0, compl1); - auto cmp1 = rewriter.create(loc, compl2, compl3); - auto cmp2 = rewriter.create(loc, cmp0, cmp1); + auto zero = mlir::arith::ConstantIndexOp::create(rewriter, loc, 0); + auto compl0 = mlir::arith::CmpIOp::create( + rewriter, loc, arith::CmpIPredicate::slt, zero, step); + auto compl1 = mlir::arith::CmpIOp::create( + rewriter, loc, arith::CmpIPredicate::sle, iv, upperBound); + auto compl2 = mlir::arith::CmpIOp::create( + rewriter, loc, arith::CmpIPredicate::slt, step, zero); + auto compl3 = mlir::arith::CmpIOp::create( + rewriter, loc, arith::CmpIPredicate::sle, upperBound, iv); + auto cmp0 = mlir::arith::AndIOp::create(rewriter, loc, compl0, compl1); + auto cmp1 = mlir::arith::AndIOp::create(rewriter, loc, compl2, compl3); + auto cmp2 = mlir::arith::OrIOp::create(rewriter, loc, cmp0, cmp1); // Remember to AND in the early-exit bool. auto comparison = - rewriter.create(loc, iterateVar, cmp2); - rewriter.create( - loc, comparison, firstBodyBlock, llvm::ArrayRef(), - endBlock, llvm::ArrayRef()); + mlir::arith::AndIOp::create(rewriter, loc, iterateVar, cmp2); + mlir::cf::CondBranchOp::create(rewriter, loc, comparison, firstBodyBlock, + llvm::ArrayRef(), endBlock, + llvm::ArrayRef()); // The result of the loop operation is the values of the condition block // arguments except the induction variable on the last iteration. auto args = whileOp.getFinalValue() diff --git a/flang/lib/Optimizer/Transforms/DebugTypeGenerator.cpp b/flang/lib/Optimizer/Transforms/DebugTypeGenerator.cpp index abad500d3f657..5dcb54eaf9b9d 100644 --- a/flang/lib/Optimizer/Transforms/DebugTypeGenerator.cpp +++ b/flang/lib/Optimizer/Transforms/DebugTypeGenerator.cpp @@ -103,13 +103,14 @@ mlir::LLVM::DILocalVariableAttr DebugTypeGenerator::generateArtificialVariable( mlir::Type type = val.getType(); if (!mlir::isa(type) || !type.isSignlessInteger()) { type = builder.getIntegerType(64); - val = builder.create(declOp.getLoc(), type, val); + val = fir::ConvertOp::create(builder, declOp.getLoc(), type, val); } mlir::LLVM::DITypeAttr Ty = convertType(type, fileAttr, scope, declOp); auto lvAttr = mlir::LLVM::DILocalVariableAttr::get( context, scope, name, fileAttr, /*line=*/0, /*argNo=*/0, /*alignInBits=*/0, Ty, mlir::LLVM::DIFlags::Artificial); - builder.create(declOp.getLoc(), val, lvAttr, nullptr); + mlir::LLVM::DbgValueOp::create(builder, declOp.getLoc(), val, lvAttr, + nullptr); return lvAttr; } diff --git a/flang/lib/Optimizer/Transforms/FIRToSCF.cpp b/flang/lib/Optimizer/Transforms/FIRToSCF.cpp index f06ad2db90d55..d7d1865bc56ba 100644 --- a/flang/lib/Optimizer/Transforms/FIRToSCF.cpp +++ b/flang/lib/Optimizer/Transforms/FIRToSCF.cpp @@ -49,13 +49,13 @@ struct DoLoopConversion : public OpRewritePattern { // must be a positive value. // For easier conversion, we calculate the trip count and use a canonical // induction variable. - auto diff = rewriter.create(loc, high, low); - auto distance = rewriter.create(loc, diff, step); - auto tripCount = rewriter.create(loc, distance, step); - auto zero = rewriter.create(loc, 0); - auto one = rewriter.create(loc, 1); + auto diff = arith::SubIOp::create(rewriter, loc, high, low); + auto distance = arith::AddIOp::create(rewriter, loc, diff, step); + auto tripCount = arith::DivSIOp::create(rewriter, loc, distance, step); + auto zero = arith::ConstantIndexOp::create(rewriter, loc, 0); + auto one = arith::ConstantIndexOp::create(rewriter, loc, 1); auto scfForOp = - rewriter.create(loc, zero, tripCount, one, iterArgs); + scf::ForOp::create(rewriter, loc, zero, tripCount, one, iterArgs); auto &loopOps = doLoopOp.getBody()->getOperations(); auto resultOp = cast(doLoopOp.getBody()->getTerminator()); @@ -68,12 +68,12 @@ struct DoLoopConversion : public OpRewritePattern { rewriter.setInsertionPointToStart(loweredBody); Value iv = - rewriter.create(loc, scfForOp.getInductionVar(), step); - iv = rewriter.create(loc, low, iv); + arith::MulIOp::create(rewriter, loc, scfForOp.getInductionVar(), step); + iv = arith::AddIOp::create(rewriter, loc, low, iv); if (!results.empty()) { rewriter.setInsertionPointToEnd(loweredBody); - rewriter.create(resultOp->getLoc(), results); + scf::YieldOp::create(rewriter, resultOp->getLoc(), results); } doLoopOp.getInductionVar().replaceAllUsesWith(iv); rewriter.replaceAllUsesWith(doLoopOp.getRegionIterArgs(), diff --git a/flang/lib/Optimizer/Transforms/GenRuntimeCallsForTest.cpp b/flang/lib/Optimizer/Transforms/GenRuntimeCallsForTest.cpp index 7ea3b9c670c69..699be12178881 100644 --- a/flang/lib/Optimizer/Transforms/GenRuntimeCallsForTest.cpp +++ b/flang/lib/Optimizer/Transforms/GenRuntimeCallsForTest.cpp @@ -91,8 +91,8 @@ void GenRuntimeCallsForTestPass::runOnOperation() { // Generate the wrapper function body that consists of a call and return. builder.setInsertionPointToStart(callerFunc.addEntryBlock()); mlir::Block::BlockArgListType args = callerFunc.front().getArguments(); - auto callOp = builder.create(loc, funcOp, args); - builder.create(loc, callOp.getResults()); + auto callOp = fir::CallOp::create(builder, loc, funcOp, args); + mlir::func::ReturnOp::create(builder, loc, callOp.getResults()); newFuncs.push_back(callerFunc.getOperation()); builder.restoreInsertionPoint(insertPt); diff --git a/flang/lib/Optimizer/Transforms/LoopVersioning.cpp b/flang/lib/Optimizer/Transforms/LoopVersioning.cpp index 056bdf63d914f..0095159398155 100644 --- a/flang/lib/Optimizer/Transforms/LoopVersioning.cpp +++ b/flang/lib/Optimizer/Transforms/LoopVersioning.cpp @@ -285,7 +285,7 @@ static mlir::Value getIndex(fir::FirOpBuilder &builder, mlir::Operation *op, // index_0 = index - lb; if (lb.getType() != index.getType()) lb = builder.createConvert(coop.getLoc(), index.getType(), lb); - return builder.create(coop.getLoc(), index, lb); + return mlir::arith::SubIOp::create(builder, coop.getLoc(), index, lb); } void LoopVersioningPass::runOnOperation() { @@ -483,26 +483,26 @@ void LoopVersioningPass::runOnOperation() { unsigned ndims = arg.rank; for (unsigned i = 0; i < ndims; i++) { mlir::Value dimIdx = builder.createIntegerConstant(loc, idxTy, i); - arg.dims[i] = builder.create(loc, idxTy, idxTy, idxTy, - arg.arg, dimIdx); + arg.dims[i] = fir::BoxDimsOp::create(builder, loc, idxTy, idxTy, idxTy, + arg.arg, dimIdx); } // We only care about lowest order dimension, here. mlir::Value elemSize = builder.createIntegerConstant(loc, idxTy, arg.size); - mlir::Value cmp = builder.create( - loc, mlir::arith::CmpIPredicate::eq, arg.dims[0].getResult(2), - elemSize); + mlir::Value cmp = mlir::arith::CmpIOp::create( + builder, loc, mlir::arith::CmpIPredicate::eq, + arg.dims[0].getResult(2), elemSize); if (!allCompares) { allCompares = cmp; } else { allCompares = - builder.create(loc, cmp, allCompares); + mlir::arith::AndIOp::create(builder, loc, cmp, allCompares); } } auto ifOp = - builder.create(loc, op.op->getResultTypes(), allCompares, - /*withElse=*/true); + fir::IfOp::create(builder, loc, op.op->getResultTypes(), allCompares, + /*withElse=*/true); builder.setInsertionPointToStart(&ifOp.getThenRegion().front()); LLVM_DEBUG(llvm::dbgs() << "Creating cloned loop\n"); @@ -515,8 +515,8 @@ void LoopVersioningPass::runOnOperation() { mlir::Type arrTy = fir::SequenceType::get(newShape, elementType); mlir::Type boxArrTy = fir::BoxType::get(arrTy); mlir::Type refArrTy = builder.getRefType(arrTy); - auto carg = builder.create(loc, boxArrTy, arg.arg); - auto caddr = builder.create(loc, refArrTy, carg); + auto carg = fir::ConvertOp::create(builder, loc, boxArrTy, arg.arg); + auto caddr = fir::BoxAddrOp::create(builder, loc, refArrTy, carg); auto insPt = builder.saveInsertionPoint(); // Use caddr instead of arg. clonedLoop->walk([&](mlir::Operation *coop) { @@ -540,9 +540,9 @@ void LoopVersioningPass::runOnOperation() { mlir::Value scale = builder.createConvert(loc, idxTy, arg.dims[i].getResult(2)); curIndex = - builder.create(loc, scale, curIndex); - totalIndex = (totalIndex) ? builder.create( - loc, curIndex, totalIndex) + mlir::arith::MulIOp::create(builder, loc, scale, curIndex); + totalIndex = (totalIndex) ? mlir::arith::AddIOp::create( + builder, loc, curIndex, totalIndex) : curIndex; } // This is the lowest dimension - which doesn't need scaling @@ -554,16 +554,16 @@ void LoopVersioningPass::runOnOperation() { unsigned bits = llvm::Log2_32(arg.size); mlir::Value elemShift = builder.createIntegerConstant(loc, idxTy, bits); - totalIndex = builder.create( - loc, - builder.create(loc, totalIndex, - elemShift), + totalIndex = mlir::arith::AddIOp::create( + builder, loc, + mlir::arith::ShRSIOp::create(builder, loc, totalIndex, + elemShift), finalIndex); } else { totalIndex = finalIndex; } - auto newOp = builder.create( - loc, builder.getRefType(elementType), caddr, + auto newOp = fir::CoordinateOp::create( + builder, loc, builder.getRefType(elementType), caddr, mlir::ValueRange{totalIndex}); LLVM_DEBUG(newOp->dump()); coop->getResult(0).replaceAllUsesWith(newOp->getResult(0)); @@ -582,7 +582,7 @@ void LoopVersioningPass::runOnOperation() { mlir::ResultRange results = clonedLoop->getResults(); bool hasResults = (results.size() > 0); if (hasResults) - builder.create(loc, results); + fir::ResultOp::create(builder, loc, results); // Add the original loop in the else-side of the if operation. builder.setInsertionPointToStart(&ifOp.getElseRegion().front()); @@ -591,7 +591,7 @@ void LoopVersioningPass::runOnOperation() { builder.insert(op.op); // Rely on "cloned loop has results, so original loop also has results". if (hasResults) { - builder.create(loc, op.op->getResults()); + fir::ResultOp::create(builder, loc, op.op->getResults()); } else { // Use an assert to check this. assert(op.op->getResults().size() == 0 && diff --git a/flang/lib/Optimizer/Transforms/MemoryAllocation.cpp b/flang/lib/Optimizer/Transforms/MemoryAllocation.cpp index 3f308a8f4b560..99040898728bb 100644 --- a/flang/lib/Optimizer/Transforms/MemoryAllocation.cpp +++ b/flang/lib/Optimizer/Transforms/MemoryAllocation.cpp @@ -68,9 +68,9 @@ static mlir::Value genAllocmem(mlir::OpBuilder &builder, fir::AllocaOp alloca, }; llvm::StringRef uniqName = unpackName(alloca.getUniqName()); llvm::StringRef bindcName = unpackName(alloca.getBindcName()); - auto heap = builder.create(alloca.getLoc(), varTy, uniqName, - bindcName, alloca.getTypeparams(), - alloca.getShape()); + auto heap = fir::AllocMemOp::create(builder, alloca.getLoc(), varTy, uniqName, + bindcName, alloca.getTypeparams(), + alloca.getShape()); LLVM_DEBUG(llvm::dbgs() << "memory allocation opt: replaced " << alloca << " with " << heap << '\n'); return heap; @@ -78,7 +78,7 @@ static mlir::Value genAllocmem(mlir::OpBuilder &builder, fir::AllocaOp alloca, static void genFreemem(mlir::Location loc, mlir::OpBuilder &builder, mlir::Value allocmem) { - [[maybe_unused]] auto free = builder.create(loc, allocmem); + [[maybe_unused]] auto free = fir::FreeMemOp::create(builder, loc, allocmem); LLVM_DEBUG(llvm::dbgs() << "memory allocation opt: add free " << free << " for " << allocmem << '\n'); } diff --git a/flang/lib/Optimizer/Transforms/MemoryUtils.cpp b/flang/lib/Optimizer/Transforms/MemoryUtils.cpp index bc4fcd8b0112e..789111cd35f67 100644 --- a/flang/lib/Optimizer/Transforms/MemoryUtils.cpp +++ b/flang/lib/Optimizer/Transforms/MemoryUtils.cpp @@ -200,30 +200,30 @@ void AllocaReplaceImpl::genIndirectDeallocation( // and access it indirectly in the entry points that do not dominate. rewriter.setInsertionPointToStart(&owningRegion.front()); mlir::Type heapType = fir::HeapType::get(alloca.getInType()); - mlir::Value ptrVar = rewriter.create(loc, heapType); - mlir::Value nullPtr = rewriter.create(loc, heapType); - rewriter.create(loc, nullPtr, ptrVar); + mlir::Value ptrVar = fir::AllocaOp::create(rewriter, loc, heapType); + mlir::Value nullPtr = fir::ZeroOp::create(rewriter, loc, heapType); + fir::StoreOp::create(rewriter, loc, nullPtr, ptrVar); // TODO: introducing a pointer compare op in FIR would help // generating less IR here. mlir::Type intPtrTy = fir::getIntPtrType(rewriter); - mlir::Value c0 = rewriter.create( - loc, intPtrTy, rewriter.getIntegerAttr(intPtrTy, 0)); + mlir::Value c0 = mlir::arith::ConstantOp::create( + rewriter, loc, intPtrTy, rewriter.getIntegerAttr(intPtrTy, 0)); // Store new storage address right after its creation. rewriter.restoreInsertionPoint(replacementInsertPoint); mlir::Value castReplacement = fir::factory::createConvert(rewriter, loc, heapType, replacement); - rewriter.create(loc, castReplacement, ptrVar); + fir::StoreOp::create(rewriter, loc, castReplacement, ptrVar); // Generate conditional deallocation at every deallocation point. auto genConditionalDealloc = [&](mlir::Location loc) { - mlir::Value ptrVal = rewriter.create(loc, ptrVar); + mlir::Value ptrVal = fir::LoadOp::create(rewriter, loc, ptrVar); mlir::Value ptrToInt = - rewriter.create(loc, intPtrTy, ptrVal); - mlir::Value isAllocated = rewriter.create( - loc, mlir::arith::CmpIPredicate::ne, ptrToInt, c0); - auto ifOp = rewriter.create(loc, mlir::TypeRange{}, isAllocated, - /*withElseRegion=*/false); + fir::ConvertOp::create(rewriter, loc, intPtrTy, ptrVal); + mlir::Value isAllocated = mlir::arith::CmpIOp::create( + rewriter, loc, mlir::arith::CmpIPredicate::ne, ptrToInt, c0); + auto ifOp = fir::IfOp::create(rewriter, loc, mlir::TypeRange{}, isAllocated, + /*withElseRegion=*/false); rewriter.setInsertionPointToStart(&ifOp.getThenRegion().front()); mlir::Value cast = fir::factory::createConvert( rewriter, loc, replacement.getType(), ptrVal); diff --git a/flang/lib/Optimizer/Transforms/PolymorphicOpConversion.cpp b/flang/lib/Optimizer/Transforms/PolymorphicOpConversion.cpp index 6e45aae4246d0..2c6601dec6e16 100644 --- a/flang/lib/Optimizer/Transforms/PolymorphicOpConversion.cpp +++ b/flang/lib/Optimizer/Transforms/PolymorphicOpConversion.cpp @@ -183,49 +183,51 @@ struct DispatchOpConv : public OpConversionPattern { mlir::Type tdescType = fir::TypeDescType::get(mlir::NoneType::get(rewriter.getContext())); mlir::Value boxDesc = - rewriter.create(loc, tdescType, passedObject); - boxDesc = rewriter.create( - loc, fir::ReferenceType::get(typeDescTy), boxDesc); + fir::BoxTypeDescOp::create(rewriter, loc, tdescType, passedObject); + boxDesc = fir::ConvertOp::create( + rewriter, loc, fir::ReferenceType::get(typeDescTy), boxDesc); // Load the bindings descriptor. auto bindingsCompName = Fortran::semantics::bindingDescCompName; fir::RecordType typeDescRecTy = mlir::cast(typeDescTy); - mlir::Value field = rewriter.create( - loc, fieldTy, bindingsCompName, typeDescRecTy, mlir::ValueRange{}); + mlir::Value field = + fir::FieldIndexOp::create(rewriter, loc, fieldTy, bindingsCompName, + typeDescRecTy, mlir::ValueRange{}); mlir::Type coorTy = fir::ReferenceType::get(typeDescRecTy.getType(bindingsCompName)); mlir::Value bindingBoxAddr = - rewriter.create(loc, coorTy, boxDesc, field); - mlir::Value bindingBox = rewriter.create(loc, bindingBoxAddr); + fir::CoordinateOp::create(rewriter, loc, coorTy, boxDesc, field); + mlir::Value bindingBox = fir::LoadOp::create(rewriter, loc, bindingBoxAddr); // Load the correct binding. - mlir::Value bindings = rewriter.create(loc, bindingBox); + mlir::Value bindings = fir::BoxAddrOp::create(rewriter, loc, bindingBox); fir::RecordType bindingTy = fir::unwrapIfDerived( mlir::cast(bindingBox.getType())); mlir::Type bindingAddrTy = fir::ReferenceType::get(bindingTy); - mlir::Value bindingIdxVal = rewriter.create( - loc, rewriter.getIndexType(), rewriter.getIndexAttr(bindingIdx)); - mlir::Value bindingAddr = rewriter.create( - loc, bindingAddrTy, bindings, bindingIdxVal); + mlir::Value bindingIdxVal = + mlir::arith::ConstantOp::create(rewriter, loc, rewriter.getIndexType(), + rewriter.getIndexAttr(bindingIdx)); + mlir::Value bindingAddr = fir::CoordinateOp::create( + rewriter, loc, bindingAddrTy, bindings, bindingIdxVal); // Get the function pointer. auto procCompName = Fortran::semantics::procCompName; - mlir::Value procField = rewriter.create( - loc, fieldTy, procCompName, bindingTy, mlir::ValueRange{}); + mlir::Value procField = fir::FieldIndexOp::create( + rewriter, loc, fieldTy, procCompName, bindingTy, mlir::ValueRange{}); fir::RecordType procTy = mlir::cast(bindingTy.getType(procCompName)); mlir::Type procRefTy = fir::ReferenceType::get(procTy); - mlir::Value procRef = rewriter.create( - loc, procRefTy, bindingAddr, procField); + mlir::Value procRef = fir::CoordinateOp::create(rewriter, loc, procRefTy, + bindingAddr, procField); auto addressFieldName = Fortran::lower::builtin::cptrFieldName; - mlir::Value addressField = rewriter.create( - loc, fieldTy, addressFieldName, procTy, mlir::ValueRange{}); + mlir::Value addressField = fir::FieldIndexOp::create( + rewriter, loc, fieldTy, addressFieldName, procTy, mlir::ValueRange{}); mlir::Type addressTy = procTy.getType(addressFieldName); mlir::Type addressRefTy = fir::ReferenceType::get(addressTy); - mlir::Value addressRef = rewriter.create( - loc, addressRefTy, procRef, addressField); - mlir::Value address = rewriter.create(loc, addressRef); + mlir::Value addressRef = fir::CoordinateOp::create( + rewriter, loc, addressRefTy, procRef, addressField); + mlir::Value address = fir::LoadOp::create(rewriter, loc, addressRef); // Get the function type. llvm::SmallVector argTypes; @@ -237,7 +239,7 @@ struct DispatchOpConv : public OpConversionPattern { mlir::Type funTy = mlir::FunctionType::get(rewriter.getContext(), argTypes, resTypes); - mlir::Value funcPtr = rewriter.create(loc, funTy, address); + mlir::Value funcPtr = fir::ConvertOp::create(rewriter, loc, funTy, address); // Make the call. llvm::SmallVector args{funcPtr}; @@ -398,12 +400,13 @@ llvm::LogicalResult SelectTypeConv::genTypeLadderStep( if (code == 0) return mlir::emitError(loc) << "type code unavailable for " << a.getType(); - mlir::Value typeCode = rewriter.create( - loc, rewriter.getI8IntegerAttr(code)); - mlir::Value selectorTypeCode = rewriter.create( - loc, rewriter.getI8Type(), selector); - cmp = rewriter.create( - loc, mlir::arith::CmpIPredicate::eq, selectorTypeCode, typeCode); + mlir::Value typeCode = mlir::arith::ConstantOp::create( + rewriter, loc, rewriter.getI8IntegerAttr(code)); + mlir::Value selectorTypeCode = fir::BoxTypeCodeOp::create( + rewriter, loc, rewriter.getI8Type(), selector); + cmp = mlir::arith::CmpIOp::create(rewriter, loc, + mlir::arith::CmpIPredicate::eq, + selectorTypeCode, typeCode); } else { // Flang inline the kind parameter in the type descriptor so we can // directly check if the type descriptor addresses are identical for @@ -418,16 +421,16 @@ llvm::LogicalResult SelectTypeConv::genTypeLadderStep( } else if (auto a = mlir::dyn_cast(attr)) { // Retrieve the type descriptor from the type guard statement record type. assert(mlir::isa(a.getType()) && "expect fir.record type"); - mlir::Value typeDescAddr = - rewriter.create(loc, mlir::TypeAttr::get(a.getType())); + mlir::Value typeDescAddr = fir::TypeDescOp::create( + rewriter, loc, mlir::TypeAttr::get(a.getType())); mlir::Type refNoneType = ReferenceType::get(rewriter.getNoneType()); mlir::Value typeDesc = - rewriter.create(loc, refNoneType, typeDescAddr); + ConvertOp::create(rewriter, loc, refNoneType, typeDescAddr); // Prepare the selector descriptor for the runtime call. mlir::Type descNoneTy = fir::BoxType::get(rewriter.getNoneType()); mlir::Value descSelector = - rewriter.create(loc, descNoneTy, selector); + ConvertOp::create(rewriter, loc, descNoneTy, selector); // Generate runtime call. llvm::StringRef fctName = RTNAME_STRING(ClassIs); @@ -455,10 +458,10 @@ llvm::LogicalResult SelectTypeConv::genTypeLadderStep( rewriter.createBlock(dest->getParent(), mlir::Region::iterator(dest)); rewriter.setInsertionPointToEnd(thisBlock); if (destOps.has_value()) - rewriter.create(loc, cmp, dest, destOps.value(), - newBlock, mlir::ValueRange{}); + mlir::cf::CondBranchOp::create(rewriter, loc, cmp, dest, destOps.value(), + newBlock, mlir::ValueRange{}); else - rewriter.create(loc, cmp, dest, newBlock); + mlir::cf::CondBranchOp::create(rewriter, loc, cmp, dest, newBlock); rewriter.setInsertionPointToEnd(newBlock); return mlir::success(); } @@ -470,16 +473,17 @@ SelectTypeConv::genTypeDescCompare(mlir::Location loc, mlir::Value selector, mlir::PatternRewriter &rewriter) const { assert(mlir::isa(ty) && "expect fir.record type"); mlir::Value typeDescAddr = - rewriter.create(loc, mlir::TypeAttr::get(ty)); - mlir::Value selectorTdescAddr = rewriter.create( - loc, typeDescAddr.getType(), selector); + fir::TypeDescOp::create(rewriter, loc, mlir::TypeAttr::get(ty)); + mlir::Value selectorTdescAddr = fir::BoxTypeDescOp::create( + rewriter, loc, typeDescAddr.getType(), selector); auto intPtrTy = rewriter.getIndexType(); auto typeDescInt = - rewriter.create(loc, intPtrTy, typeDescAddr); + fir::ConvertOp::create(rewriter, loc, intPtrTy, typeDescAddr); auto selectorTdescInt = - rewriter.create(loc, intPtrTy, selectorTdescAddr); - return rewriter.create( - loc, mlir::arith::CmpIPredicate::eq, typeDescInt, selectorTdescInt); + fir::ConvertOp::create(rewriter, loc, intPtrTy, selectorTdescAddr); + return mlir::arith::CmpIOp::create(rewriter, loc, + mlir::arith::CmpIPredicate::eq, + typeDescInt, selectorTdescInt); } llvm::SmallSet diff --git a/flang/lib/Optimizer/Transforms/SimplifyFIROperations.cpp b/flang/lib/Optimizer/Transforms/SimplifyFIROperations.cpp index ad8464b495888..c6aec96ceb5ae 100644 --- a/flang/lib/Optimizer/Transforms/SimplifyFIROperations.cpp +++ b/flang/lib/Optimizer/Transforms/SimplifyFIROperations.cpp @@ -88,18 +88,18 @@ mlir::LogicalResult IsContiguousBoxCoversion::matchAndRewrite( // The scalar cases are supposed to be optimized by the canonicalization. if (rank == 1 || (op.getInnermost() && rank > 0)) { mlir::Type idxTy = builder.getIndexType(); - auto eleSize = builder.create(loc, idxTy, box); + auto eleSize = fir::BoxEleSizeOp::create(builder, loc, idxTy, box); mlir::Value zero = fir::factory::createZeroValue(builder, loc, idxTy); auto dimInfo = - builder.create(loc, idxTy, idxTy, idxTy, box, zero); + fir::BoxDimsOp::create(builder, loc, idxTy, idxTy, idxTy, box, zero); mlir::Value stride = dimInfo.getByteStride(); - mlir::Value pred1 = builder.create( - loc, mlir::arith::CmpIPredicate::eq, eleSize, stride); + mlir::Value pred1 = mlir::arith::CmpIOp::create( + builder, loc, mlir::arith::CmpIPredicate::eq, eleSize, stride); mlir::Value extent = dimInfo.getExtent(); - mlir::Value pred2 = builder.create( - loc, mlir::arith::CmpIPredicate::eq, extent, zero); + mlir::Value pred2 = mlir::arith::CmpIOp::create( + builder, loc, mlir::arith::CmpIPredicate::eq, extent, zero); mlir::Value result = - builder.create(loc, pred1, pred2); + mlir::arith::OrIOp::create(builder, loc, pred1, pred2); result = builder.createConvert(loc, op.getType(), result); rewriter.replaceOp(op, result); return mlir::success(); @@ -192,7 +192,7 @@ class DoConcurrentConversion // TODO Should this be a heap allocation instead? For now, we allocate // on the stack for each loop iteration. mlir::Value localAlloc = - rewriter.create(loop.getLoc(), localizer.getType()); + fir::AllocaOp::create(rewriter, loop.getLoc(), localizer.getType()); auto cloneLocalizerRegion = [&](mlir::Region ®ion, mlir::ValueRange regionArgs, @@ -258,8 +258,8 @@ class DoConcurrentConversion for (auto [lb, ub, st, iv] : llvm::zip_equal(loop.getLowerBound(), loop.getUpperBound(), loop.getStep(), *loop.getLoopInductionVars())) { - innermostUnorderdLoop = rewriter.create( - doConcurentOp.getLoc(), lb, ub, st, + innermostUnorderdLoop = fir::DoLoopOp::create( + rewriter, doConcurentOp.getLoc(), lb, ub, st, /*unordred=*/true, /*finalCountValue=*/false, /*iterArgs=*/mlir::ValueRange{}, loop.getReduceVars(), loop.getReduceAttrsAttr()); diff --git a/flang/lib/Optimizer/Transforms/SimplifyIntrinsics.cpp b/flang/lib/Optimizer/Transforms/SimplifyIntrinsics.cpp index 4d25a02bf18ba..49a085ee3b336 100644 --- a/flang/lib/Optimizer/Transforms/SimplifyIntrinsics.cpp +++ b/flang/lib/Optimizer/Transforms/SimplifyIntrinsics.cpp @@ -284,7 +284,7 @@ genReductionLoop(fir::FirOpBuilder &builder, mlir::func::FuncOp &funcOp, fir::SequenceType::getUnknownExtent()); mlir::Type arrTy = fir::SequenceType::get(flatShape, elementType); mlir::Type boxArrTy = fir::BoxType::get(arrTy); - mlir::Value array = builder.create(loc, boxArrTy, arg); + mlir::Value array = fir::ConvertOp::create(builder, loc, boxArrTy, arg); mlir::Type resultType = funcOp.getResultTypes()[0]; mlir::Value init = initVal(builder, loc, resultType); @@ -299,11 +299,11 @@ genReductionLoop(fir::FirOpBuilder &builder, mlir::func::FuncOp &funcOp, // should be able to optimize the redundancy. for (unsigned i = 0; i < rank; ++i) { mlir::Value dimIdx = builder.createIntegerConstant(loc, idxTy, i); - auto dims = - builder.create(loc, idxTy, idxTy, idxTy, array, dimIdx); + auto dims = fir::BoxDimsOp::create(builder, loc, idxTy, idxTy, idxTy, array, + dimIdx); mlir::Value len = dims.getResult(1); // We use C indexing here, so len-1 as loopcount - mlir::Value loopCount = builder.create(loc, len, one); + mlir::Value loopCount = mlir::arith::SubIOp::create(builder, loc, len, one); bounds.push_back(loopCount); } // Create a loop nest consisting of OP operations. @@ -316,9 +316,9 @@ genReductionLoop(fir::FirOpBuilder &builder, mlir::func::FuncOp &funcOp, for (unsigned i = rank; 0 < i; --i) { mlir::Value step = one; mlir::Value loopCount = bounds[i - 1]; - auto loop = builder.create(loc, zeroIdx, loopCount, step, - unorderedOrInitialLoopCond, - /*finalCountValue=*/false, init); + auto loop = OP::create(builder, loc, zeroIdx, loopCount, step, + unorderedOrInitialLoopCond, + /*finalCountValue=*/false, init); init = loop.getRegionIterArgs()[resultIndex]; indices.push_back(loop.getInductionVar()); // Set insertion point to the loop body so that the next loop @@ -332,8 +332,8 @@ genReductionLoop(fir::FirOpBuilder &builder, mlir::func::FuncOp &funcOp, // We are in the innermost loop: generate the reduction body. mlir::Type eleRefTy = builder.getRefType(elementType); mlir::Value addr = - builder.create(loc, eleRefTy, array, indices); - mlir::Value elem = builder.create(loc, addr); + fir::CoordinateOp::create(builder, loc, eleRefTy, array, indices); + mlir::Value elem = fir::LoadOp::create(builder, loc, addr); mlir::Value reductionVal = genBody(builder, loc, elementType, elem, init); // Generate vector with condition to continue while loop at [0] and result // from current loop at [1] for IterWhileOp loops, just result at [0] for @@ -344,7 +344,7 @@ genReductionLoop(fir::FirOpBuilder &builder, mlir::func::FuncOp &funcOp, // to return the updated value of the reduction to the enclosing // loops. for (unsigned i = 0; i < rank; ++i) { - auto result = builder.create(loc, results); + auto result = fir::ResultOp::create(builder, loc, results); // Proceed to the outer loop. auto loop = mlir::cast(result->getParentOp()); results = loop.getResults(); @@ -354,7 +354,7 @@ genReductionLoop(fir::FirOpBuilder &builder, mlir::func::FuncOp &funcOp, } // End of loop nest. The insertion point is after the outermost loop. // Return the reduction value from the function. - builder.create(loc, results[resultIndex]); + mlir::func::ReturnOp::create(builder, loc, results[resultIndex]); } static llvm::SmallVector nopLoopCond(fir::FirOpBuilder &builder, @@ -394,9 +394,9 @@ static void genRuntimeSumBody(fir::FirOpBuilder &builder, mlir::Type elementType, mlir::Value elem1, mlir::Value elem2) -> mlir::Value { if (mlir::isa(elementType)) - return builder.create(loc, elem1, elem2); + return mlir::arith::AddFOp::create(builder, loc, elem1, elem2); if (mlir::isa(elementType)) - return builder.create(loc, elem1, elem2); + return mlir::arith::AddIOp::create(builder, loc, elem1, elem2); llvm_unreachable("unsupported type"); return {}; @@ -436,12 +436,12 @@ static void genRuntimeMaxvalBody(fir::FirOpBuilder &builder, // This libm function may not work properly for F128 arguments // on targets where long double is not F128. It is an LLVM issue, // but we just use normal select here to resolve all the cases. - auto compare = builder.create( - loc, mlir::arith::CmpFPredicate::OGT, elem1, elem2); - return builder.create(loc, compare, elem1, elem2); + auto compare = mlir::arith::CmpFOp::create( + builder, loc, mlir::arith::CmpFPredicate::OGT, elem1, elem2); + return mlir::arith::SelectOp::create(builder, loc, compare, elem1, elem2); } if (mlir::isa(elementType)) - return builder.create(loc, elem1, elem2); + return mlir::arith::MaxSIOp::create(builder, loc, elem1, elem2); llvm_unreachable("unsupported type"); return {}; @@ -472,11 +472,11 @@ static void genRuntimeCountBody(fir::FirOpBuilder &builder, auto zero64 = builder.createIntegerConstant(loc, builder.getI64Type(), 0); auto one64 = builder.createIntegerConstant(loc, builder.getI64Type(), 1); - auto compare = builder.create( - loc, mlir::arith::CmpIPredicate::eq, elem1, zero32); + auto compare = mlir::arith::CmpIOp::create( + builder, loc, mlir::arith::CmpIPredicate::eq, elem1, zero32); auto select = - builder.create(loc, compare, zero64, one64); - return builder.create(loc, select, elem2); + mlir::arith::SelectOp::create(builder, loc, compare, zero64, one64); + return mlir::arith::AddIOp::create(builder, loc, select, elem2); }; // Count always gets I32 for elementType as it converts logical input to @@ -501,14 +501,14 @@ static void genRuntimeAnyBody(fir::FirOpBuilder &builder, mlir::Type elementType, mlir::Value elem1, mlir::Value elem2) -> mlir::Value { auto zero = builder.createIntegerConstant(loc, elementType, 0); - return builder.create( - loc, mlir::arith::CmpIPredicate::ne, elem1, zero); + return mlir::arith::CmpIOp::create( + builder, loc, mlir::arith::CmpIPredicate::ne, elem1, zero); }; auto continueCond = [](fir::FirOpBuilder builder, mlir::Location loc, mlir::Value reductionVal) { auto one1 = builder.createIntegerConstant(loc, builder.getI1Type(), 1); - auto eor = builder.create(loc, reductionVal, one1); + auto eor = mlir::arith::XOrIOp::create(builder, loc, reductionVal, one1); llvm::SmallVector results = {eor, reductionVal}; return results; }; @@ -534,8 +534,8 @@ static void genRuntimeAllBody(fir::FirOpBuilder &builder, mlir::Type elementType, mlir::Value elem1, mlir::Value elem2) -> mlir::Value { auto zero = builder.createIntegerConstant(loc, elementType, 0); - return builder.create( - loc, mlir::arith::CmpIPredicate::ne, elem1, zero); + return mlir::arith::CmpIOp::create( + builder, loc, mlir::arith::CmpIPredicate::ne, elem1, zero); }; auto continueCond = [](fir::FirOpBuilder builder, mlir::Location loc, @@ -577,13 +577,13 @@ void fir::genMinMaxlocReductionLoop( fir::SequenceType::getUnknownExtent()); mlir::Type arrTy = fir::SequenceType::get(flatShape, elementType); mlir::Type boxArrTy = fir::BoxType::get(arrTy); - array = builder.create(loc, boxArrTy, array); + array = fir::ConvertOp::create(builder, loc, boxArrTy, array); mlir::Type resultElemType = hlfir::getFortranElementType(resultArr.getType()); mlir::Value flagSet = builder.createIntegerConstant(loc, resultElemType, 1); mlir::Value zero = builder.createIntegerConstant(loc, resultElemType, 0); mlir::Value flagRef = builder.createTemporary(loc, resultElemType); - builder.create(loc, zero, flagRef); + fir::StoreOp::create(builder, loc, zero, flagRef); mlir::Value init = initVal(builder, loc, elementType); llvm::SmallVector bounds; @@ -597,11 +597,11 @@ void fir::genMinMaxlocReductionLoop( // should be able to optimize the redundancy. for (unsigned i = 0; i < rank; ++i) { mlir::Value dimIdx = builder.createIntegerConstant(loc, idxTy, i); - auto dims = - builder.create(loc, idxTy, idxTy, idxTy, array, dimIdx); + auto dims = fir::BoxDimsOp::create(builder, loc, idxTy, idxTy, idxTy, array, + dimIdx); mlir::Value len = dims.getResult(1); // We use C indexing here, so len-1 as loopcount - mlir::Value loopCount = builder.create(loc, len, one); + mlir::Value loopCount = mlir::arith::SubIOp::create(builder, loc, len, one); bounds.push_back(loopCount); } // Create a loop nest consisting of OP operations. @@ -615,8 +615,8 @@ void fir::genMinMaxlocReductionLoop( mlir::Value step = one; mlir::Value loopCount = bounds[i - 1]; auto loop = - builder.create(loc, zeroIdx, loopCount, step, false, - /*finalCountValue=*/false, init); + fir::DoLoopOp::create(builder, loc, zeroIdx, loopCount, step, false, + /*finalCountValue=*/false, init); init = loop.getRegionIterArgs()[0]; indices.push_back(loop.getInductionVar()); // Set insertion point to the loop body so that the next loop @@ -634,7 +634,7 @@ void fir::genMinMaxlocReductionLoop( // to return the updated value of the reduction to the enclosing // loops. for (unsigned i = 0; i < rank; ++i) { - auto result = builder.create(loc, reductionVal); + auto result = fir::ResultOp::create(builder, loc, reductionVal); // Proceed to the outer loop. auto loop = mlir::cast(result->getParentOp()); reductionVal = loop.getResult(0); @@ -646,7 +646,7 @@ void fir::genMinMaxlocReductionLoop( if (maskMayBeLogicalScalar) { if (fir::IfOp ifOp = mlir::dyn_cast(builder.getBlock()->getParentOp())) { - builder.create(loc, reductionVal); + fir::ResultOp::create(builder, loc, reductionVal); builder.setInsertionPointAfter(ifOp); // Redefine flagSet to escape scope of ifOp flagSet = builder.createIntegerConstant(loc, resultElemType, 1); @@ -689,10 +689,11 @@ static void genRuntimeMinMaxlocBody(fir::FirOpBuilder &builder, mlir::Value returnValue = builder.createIntegerConstant(loc, resultElemTy, 0); mlir::Value resultArrSize = builder.createIntegerConstant(loc, idxTy, rank); - mlir::Value resultArrInit = builder.create(loc, resultTy); - mlir::Value resultArrShape = builder.create(loc, resultArrSize); - mlir::Value resultArr = builder.create( - loc, resultBoxTy, resultArrInit, resultArrShape); + mlir::Value resultArrInit = fir::AllocMemOp::create(builder, loc, resultTy); + mlir::Value resultArrShape = + fir::ShapeOp::create(builder, loc, resultArrSize); + mlir::Value resultArr = fir::EmboxOp::create(builder, loc, resultBoxTy, + resultArrInit, resultArrShape); mlir::Type resultRefTy = builder.getRefType(resultElemTy); @@ -701,14 +702,14 @@ static void genRuntimeMinMaxlocBody(fir::FirOpBuilder &builder, fir::SequenceType::getUnknownExtent()); mlir::Type maskTy = fir::SequenceType::get(flatShape, maskElemType); mlir::Type boxMaskTy = fir::BoxType::get(maskTy); - mask = builder.create(loc, boxMaskTy, mask); + mask = fir::ConvertOp::create(builder, loc, boxMaskTy, mask); } for (unsigned int i = 0; i < rank; ++i) { mlir::Value index = builder.createIntegerConstant(loc, idxTy, i); mlir::Value resultElemAddr = - builder.create(loc, resultRefTy, resultArr, index); - builder.create(loc, returnValue, resultElemAddr); + fir::CoordinateOp::create(builder, loc, resultRefTy, resultArr, index); + fir::StoreOp::create(builder, loc, returnValue, resultElemAddr); } auto genBodyOp = @@ -720,29 +721,30 @@ static void genRuntimeMinMaxlocBody(fir::FirOpBuilder &builder, if (maskRank > 0) { mlir::Type logicalRef = builder.getRefType(maskElemType); mlir::Value maskAddr = - builder.create(loc, logicalRef, mask, indices); - mlir::Value maskElem = builder.create(loc, maskAddr); + fir::CoordinateOp::create(builder, loc, logicalRef, mask, indices); + mlir::Value maskElem = fir::LoadOp::create(builder, loc, maskAddr); // fir::IfOp requires argument to be I1 - won't accept logical or any // other Integer. mlir::Type ifCompatType = builder.getI1Type(); mlir::Value ifCompatElem = - builder.create(loc, ifCompatType, maskElem); + fir::ConvertOp::create(builder, loc, ifCompatType, maskElem); llvm::SmallVector resultsTy = {elementType, elementType}; - fir::IfOp ifOp = builder.create(loc, elementType, ifCompatElem, - /*withElseRegion=*/true); + fir::IfOp ifOp = + fir::IfOp::create(builder, loc, elementType, ifCompatElem, + /*withElseRegion=*/true); builder.setInsertionPointToStart(&ifOp.getThenRegion().front()); } // Set flag that mask was true at some point mlir::Value flagSet = builder.createIntegerConstant( loc, mlir::cast(flagRef.getType()).getEleTy(), 1); - mlir::Value isFirst = builder.create(loc, flagRef); + mlir::Value isFirst = fir::LoadOp::create(builder, loc, flagRef); mlir::Type eleRefTy = builder.getRefType(elementType); mlir::Value addr = - builder.create(loc, eleRefTy, array, indices); - mlir::Value elem = builder.create(loc, addr); + fir::CoordinateOp::create(builder, loc, eleRefTy, array, indices); + mlir::Value elem = fir::LoadOp::create(builder, loc, addr); mlir::Value cmp; if (mlir::isa(elementType)) { @@ -750,38 +752,37 @@ static void genRuntimeMinMaxlocBody(fir::FirOpBuilder &builder, // is not NaN. A OGL/OLT condition will usually work for this unless all // the values are Nan or Inf. This follows the same logic as // NumericCompare for Minloc/Maxlox in extrema.cpp. - cmp = builder.create( - loc, - isMax ? mlir::arith::CmpFPredicate::OGT - : mlir::arith::CmpFPredicate::OLT, - elem, reduction); - - mlir::Value cmpNan = builder.create( - loc, mlir::arith::CmpFPredicate::UNE, reduction, reduction); - mlir::Value cmpNan2 = builder.create( - loc, mlir::arith::CmpFPredicate::OEQ, elem, elem); - cmpNan = builder.create(loc, cmpNan, cmpNan2); - cmp = builder.create(loc, cmp, cmpNan); + cmp = mlir::arith::CmpFOp::create(builder, loc, + isMax ? mlir::arith::CmpFPredicate::OGT + : mlir::arith::CmpFPredicate::OLT, + elem, reduction); + + mlir::Value cmpNan = mlir::arith::CmpFOp::create( + builder, loc, mlir::arith::CmpFPredicate::UNE, reduction, reduction); + mlir::Value cmpNan2 = mlir::arith::CmpFOp::create( + builder, loc, mlir::arith::CmpFPredicate::OEQ, elem, elem); + cmpNan = mlir::arith::AndIOp::create(builder, loc, cmpNan, cmpNan2); + cmp = mlir::arith::OrIOp::create(builder, loc, cmp, cmpNan); } else if (mlir::isa(elementType)) { - cmp = builder.create( - loc, - isMax ? mlir::arith::CmpIPredicate::sgt - : mlir::arith::CmpIPredicate::slt, - elem, reduction); + cmp = mlir::arith::CmpIOp::create(builder, loc, + isMax ? mlir::arith::CmpIPredicate::sgt + : mlir::arith::CmpIPredicate::slt, + elem, reduction); } else { llvm_unreachable("unsupported type"); } // The condition used for the loop is isFirst || . - isFirst = builder.create(loc, cmp.getType(), isFirst); - isFirst = builder.create( - loc, isFirst, builder.createIntegerConstant(loc, cmp.getType(), 1)); - cmp = builder.create(loc, cmp, isFirst); - fir::IfOp ifOp = builder.create(loc, elementType, cmp, - /*withElseRegion*/ true); + isFirst = fir::ConvertOp::create(builder, loc, cmp.getType(), isFirst); + isFirst = mlir::arith::XOrIOp::create( + builder, loc, isFirst, + builder.createIntegerConstant(loc, cmp.getType(), 1)); + cmp = mlir::arith::OrIOp::create(builder, loc, cmp, isFirst); + fir::IfOp ifOp = fir::IfOp::create(builder, loc, elementType, cmp, + /*withElseRegion*/ true); builder.setInsertionPointToStart(&ifOp.getThenRegion().front()); - builder.create(loc, flagSet, flagRef); + fir::StoreOp::create(builder, loc, flagSet, flagRef); mlir::Type resultElemTy = hlfir::getFortranElementType(resultArr.getType()); mlir::Type returnRefTy = builder.getRefType(resultElemTy); mlir::IndexType idxTy = builder.getIndexType(); @@ -790,17 +791,17 @@ static void genRuntimeMinMaxlocBody(fir::FirOpBuilder &builder, for (unsigned int i = 0; i < rank; ++i) { mlir::Value index = builder.createIntegerConstant(loc, idxTy, i); - mlir::Value resultElemAddr = - builder.create(loc, returnRefTy, resultArr, index); + mlir::Value resultElemAddr = fir::CoordinateOp::create( + builder, loc, returnRefTy, resultArr, index); mlir::Value convert = - builder.create(loc, resultElemTy, indices[i]); + fir::ConvertOp::create(builder, loc, resultElemTy, indices[i]); mlir::Value fortranIndex = - builder.create(loc, convert, one); - builder.create(loc, fortranIndex, resultElemAddr); + mlir::arith::AddIOp::create(builder, loc, convert, one); + fir::StoreOp::create(builder, loc, fortranIndex, resultElemAddr); } - builder.create(loc, elem); + fir::ResultOp::create(builder, loc, elem); builder.setInsertionPointToStart(&ifOp.getElseRegion().front()); - builder.create(loc, reduction); + fir::ResultOp::create(builder, loc, reduction); builder.setInsertionPointAfter(ifOp); mlir::Value reductionVal = ifOp.getResult(0); @@ -808,9 +809,9 @@ static void genRuntimeMinMaxlocBody(fir::FirOpBuilder &builder, if (maskRank > 0) { fir::IfOp ifOp = mlir::dyn_cast(builder.getBlock()->getParentOp()); - builder.create(loc, reductionVal); + fir::ResultOp::create(builder, loc, reductionVal); builder.setInsertionPointToStart(&ifOp.getElseRegion().front()); - builder.create(loc, reduction); + fir::ResultOp::create(builder, loc, reduction); reductionVal = ifOp.getResult(0); builder.setInsertionPointAfter(ifOp); } @@ -825,12 +826,12 @@ static void genRuntimeMinMaxlocBody(fir::FirOpBuilder &builder, mlir::Type logical = maskElemType; mlir::Type logicalRefTy = builder.getRefType(logical); mlir::Value condAddr = - builder.create(loc, logicalRefTy, mask); - mlir::Value cond = builder.create(loc, condAddr); - mlir::Value condI1 = builder.create(loc, i1Type, cond); + fir::BoxAddrOp::create(builder, loc, logicalRefTy, mask); + mlir::Value cond = fir::LoadOp::create(builder, loc, condAddr); + mlir::Value condI1 = fir::ConvertOp::create(builder, loc, i1Type, cond); - fir::IfOp ifOp = builder.create(loc, elementType, condI1, - /*withElseRegion=*/true); + fir::IfOp ifOp = fir::IfOp::create(builder, loc, elementType, condI1, + /*withElseRegion=*/true); builder.setInsertionPointToStart(&ifOp.getElseRegion().front()); mlir::Value basicValue; @@ -839,7 +840,7 @@ static void genRuntimeMinMaxlocBody(fir::FirOpBuilder &builder, } else { basicValue = builder.createRealConstant(loc, elementType, 0); } - builder.create(loc, basicValue); + fir::ResultOp::create(builder, loc, basicValue); builder.setInsertionPointToStart(&ifOp.getThenRegion().front()); } @@ -847,8 +848,8 @@ static void genRuntimeMinMaxlocBody(fir::FirOpBuilder &builder, const mlir::Type &resultElemType, mlir::Value resultArr, mlir::Value index) { mlir::Type resultRefTy = builder.getRefType(resultElemType); - return builder.create(loc, resultRefTy, resultArr, - index); + return fir::CoordinateOp::create(builder, loc, resultRefTy, resultArr, + index); }; genMinMaxlocReductionLoop(builder, funcOp.front().getArgument(1), init, @@ -859,25 +860,26 @@ static void genRuntimeMinMaxlocBody(fir::FirOpBuilder &builder, if (isDim) { mlir::Type resultBoxTy = fir::BoxType::get(fir::HeapType::get(resultElemTy)); - mlir::Value outputArr = builder.create( - loc, builder.getRefType(resultBoxTy), funcOp.front().getArgument(0)); - mlir::Value resultArrScalar = builder.create( - loc, fir::HeapType::get(resultElemTy), resultArrInit); + mlir::Value outputArr = + fir::ConvertOp::create(builder, loc, builder.getRefType(resultBoxTy), + funcOp.front().getArgument(0)); + mlir::Value resultArrScalar = fir::ConvertOp::create( + builder, loc, fir::HeapType::get(resultElemTy), resultArrInit); mlir::Value resultBox = - builder.create(loc, resultBoxTy, resultArrScalar); - builder.create(loc, resultBox, outputArr); + fir::EmboxOp::create(builder, loc, resultBoxTy, resultArrScalar); + fir::StoreOp::create(builder, loc, resultBox, outputArr); } else { fir::SequenceType::Shape resultShape(1, rank); mlir::Type outputArrTy = fir::SequenceType::get(resultShape, resultElemTy); mlir::Type outputHeapTy = fir::HeapType::get(outputArrTy); mlir::Type outputBoxTy = fir::BoxType::get(outputHeapTy); mlir::Type outputRefTy = builder.getRefType(outputBoxTy); - mlir::Value outputArr = builder.create( - loc, outputRefTy, funcOp.front().getArgument(0)); - builder.create(loc, resultArr, outputArr); + mlir::Value outputArr = fir::ConvertOp::create( + builder, loc, outputRefTy, funcOp.front().getArgument(0)); + fir::StoreOp::create(builder, loc, resultArr, outputArr); } - builder.create(loc); + mlir::func::ReturnOp::create(builder, loc); } /// Generate function type for the simplified version of RTNAME(DotProduct) @@ -929,10 +931,10 @@ static void genRuntimeDotBody(fir::FirOpBuilder &builder, fir::SequenceType::Shape flatShape = {fir::SequenceType::getUnknownExtent()}; mlir::Type arrTy1 = fir::SequenceType::get(flatShape, arg1ElementTy); mlir::Type boxArrTy1 = fir::BoxType::get(arrTy1); - mlir::Value array1 = builder.create(loc, boxArrTy1, arg1); + mlir::Value array1 = fir::ConvertOp::create(builder, loc, boxArrTy1, arg1); mlir::Type arrTy2 = fir::SequenceType::get(flatShape, arg2ElementTy); mlir::Type boxArrTy2 = fir::BoxType::get(arrTy2); - mlir::Value array2 = builder.create(loc, boxArrTy2, arg2); + mlir::Value array2 = fir::ConvertOp::create(builder, loc, boxArrTy2, arg2); // This version takes the loop trip count from the first argument. // If the first argument's box has unknown (at compilation time) // extent, then it may be better to take the extent from the second @@ -941,17 +943,17 @@ static void genRuntimeDotBody(fir::FirOpBuilder &builder, // function and some analysis at the call site to choose which version // is more profitable to call. // Note that we can assume that both arguments have the same extent. - auto dims = - builder.create(loc, idxTy, idxTy, idxTy, array1, zeroIdx); + auto dims = fir::BoxDimsOp::create(builder, loc, idxTy, idxTy, idxTy, array1, + zeroIdx); mlir::Value len = dims.getResult(1); mlir::Value one = builder.createIntegerConstant(loc, idxTy, 1); mlir::Value step = one; // We use C indexing here, so len-1 as loopcount - mlir::Value loopCount = builder.create(loc, len, one); - auto loop = builder.create(loc, zeroIdx, loopCount, step, - /*unordered=*/false, - /*finalCountValue=*/false, zero); + mlir::Value loopCount = mlir::arith::SubIOp::create(builder, loc, len, one); + auto loop = fir::DoLoopOp::create(builder, loc, zeroIdx, loopCount, step, + /*unordered=*/false, + /*finalCountValue=*/false, zero); mlir::Value sumVal = loop.getRegionIterArgs()[0]; // Begin loop code @@ -961,33 +963,35 @@ static void genRuntimeDotBody(fir::FirOpBuilder &builder, mlir::Type eleRef1Ty = builder.getRefType(arg1ElementTy); mlir::Value index = loop.getInductionVar(); mlir::Value addr1 = - builder.create(loc, eleRef1Ty, array1, index); - mlir::Value elem1 = builder.create(loc, addr1); + fir::CoordinateOp::create(builder, loc, eleRef1Ty, array1, index); + mlir::Value elem1 = fir::LoadOp::create(builder, loc, addr1); // Convert to the result type. - elem1 = builder.create(loc, resultElementType, elem1); + elem1 = fir::ConvertOp::create(builder, loc, resultElementType, elem1); mlir::Type eleRef2Ty = builder.getRefType(arg2ElementTy); mlir::Value addr2 = - builder.create(loc, eleRef2Ty, array2, index); - mlir::Value elem2 = builder.create(loc, addr2); + fir::CoordinateOp::create(builder, loc, eleRef2Ty, array2, index); + mlir::Value elem2 = fir::LoadOp::create(builder, loc, addr2); // Convert to the result type. - elem2 = builder.create(loc, resultElementType, elem2); + elem2 = fir::ConvertOp::create(builder, loc, resultElementType, elem2); if (mlir::isa(resultElementType)) - sumVal = builder.create( - loc, builder.create(loc, elem1, elem2), sumVal); + sumVal = mlir::arith::AddFOp::create( + builder, loc, mlir::arith::MulFOp::create(builder, loc, elem1, elem2), + sumVal); else if (mlir::isa(resultElementType)) - sumVal = builder.create( - loc, builder.create(loc, elem1, elem2), sumVal); + sumVal = mlir::arith::AddIOp::create( + builder, loc, mlir::arith::MulIOp::create(builder, loc, elem1, elem2), + sumVal); else llvm_unreachable("unsupported type"); - builder.create(loc, sumVal); + fir::ResultOp::create(builder, loc, sumVal); // End of loop. builder.restoreInsertionPoint(loopEndPt); mlir::Value resultVal = loop.getResult(0); - builder.create(loc, resultVal); + mlir::func::ReturnOp::create(builder, loc, resultVal); } mlir::func::FuncOp SimplifyIntrinsicsPass::getOrCreateFunction( @@ -1229,8 +1233,8 @@ void SimplifyIntrinsicsPass::simplifyMinMaxlocReduction( mlir::func::FuncOp newFunc = getOrCreateFunction(builder, funcName, typeGenerator, bodyGenerator); - builder.create(loc, newFunc, - mlir::ValueRange{args[0], args[1], mask}); + fir::CallOp::create(builder, loc, newFunc, + mlir::ValueRange{args[0], args[1], mask}); call->dropAllReferences(); call->erase(); } @@ -1259,7 +1263,7 @@ void SimplifyIntrinsicsPass::simplifyReductionBody( mlir::func::FuncOp newFunc = getOrCreateFunction(builder, funcName, typeGenerator, bodyGenerator); auto newCall = - builder.create(loc, newFunc, mlir::ValueRange{args[0]}); + fir::CallOp::create(builder, loc, newFunc, mlir::ValueRange{args[0]}); call->replaceAllUsesWith(newCall.getResults()); call->dropAllReferences(); call->erase(); @@ -1344,8 +1348,8 @@ void SimplifyIntrinsicsPass::runOnOperation() { mlir::func::FuncOp newFunc = getOrCreateFunction( builder, typedFuncName, typeGenerator, bodyGenerator); - auto newCall = builder.create(loc, newFunc, - mlir::ValueRange{v1, v2}); + auto newCall = fir::CallOp::create(builder, loc, newFunc, + mlir::ValueRange{v1, v2}); call->replaceAllUsesWith(newCall.getResults()); call->dropAllReferences(); call->erase(); diff --git a/flang/lib/Optimizer/Transforms/StackArrays.cpp b/flang/lib/Optimizer/Transforms/StackArrays.cpp index bc8a9497fbb70..0d131291feef3 100644 --- a/flang/lib/Optimizer/Transforms/StackArrays.cpp +++ b/flang/lib/Optimizer/Transforms/StackArrays.cpp @@ -569,7 +569,7 @@ static mlir::Value convertAllocationType(mlir::PatternRewriter &rewriter, auto insertionPoint = rewriter.saveInsertionPoint(); rewriter.setInsertionPointAfter(stack.getDefiningOp()); mlir::Value conv = - rewriter.create(loc, firHeapTy, stack).getResult(); + fir::ConvertOp::create(rewriter, loc, firHeapTy, stack).getResult(); rewriter.restoreInsertionPoint(insertionPoint); return conv; } @@ -758,9 +758,9 @@ AllocMemConversion::insertAlloca(fir::AllocMemOp &oldAlloc, llvm::StringRef uniqName = unpackName(oldAlloc.getUniqName()); llvm::StringRef bindcName = unpackName(oldAlloc.getBindcName()); - auto alloca = rewriter.create(loc, varTy, uniqName, bindcName, - oldAlloc.getTypeparams(), - oldAlloc.getShape()); + auto alloca = + fir::AllocaOp::create(rewriter, loc, varTy, uniqName, bindcName, + oldAlloc.getTypeparams(), oldAlloc.getShape()); if (emitLifetimeMarkers) insertLifetimeMarkers(oldAlloc, alloca, rewriter);