diff --git a/include/polygeist/Dialect.td b/include/polygeist/Dialect.td index 10f3a16197cf..4d4f040c3c9d 100644 --- a/include/polygeist/Dialect.td +++ b/include/polygeist/Dialect.td @@ -25,7 +25,7 @@ def Polygeist_Dialect : Dialect { // Base BFV operation definition. //===----------------------------------------------------------------------===// -class Polygeist_Op traits = []> +class Polygeist_Op traits = []> : Op; #endif // POLYGEIST_DIALECT diff --git a/include/polygeist/Passes/Passes.td b/include/polygeist/Passes/Passes.td index 86eed651d7d6..965046c1f84a 100644 --- a/include/polygeist/Passes/Passes.td +++ b/include/polygeist/Passes/Passes.td @@ -3,12 +3,12 @@ include "mlir/Pass/PassBase.td" -def AffineCFG : FunctionPass<"affine-cfg"> { +def AffineCFG : Pass<"affine-cfg", "FuncOp"> { let summary = "Replace scf.if and similar with affine.if"; let constructor = "mlir::polygeist::replaceAffineCFGPass()"; } -def Mem2Reg : FunctionPass<"mem2reg"> { +def Mem2Reg : Pass<"mem2reg", "FuncOp"> { let summary = "Replace scf.if and similar with affine.if"; let constructor = "mlir::polygeist::createMem2RegPass()"; } @@ -20,12 +20,12 @@ def ParallelLower : Pass<"parallel-lower", "mlir::ModuleOp"> { let constructor = "mlir::polygeist::createParallelLowerPass()"; } -def AffineReduction : FunctionPass<"detect-reduction"> { +def AffineReduction : Pass<"detect-reduction", "FuncOp"> { let summary = "Detect reductions in affine.for"; let constructor = "mlir::polygeist::detectReductionPass()"; } -def SCFCPUify : FunctionPass<"cpuify"> { +def SCFCPUify : Pass<"cpuify", "FuncOp"> { let summary = "remove scf.barrier"; let constructor = "mlir::polygeist::createCPUifyPass()"; let dependentDialects = @@ -35,29 +35,29 @@ def SCFCPUify : FunctionPass<"cpuify"> { ]; } -def SCFBarrierRemovalContinuation : FunctionPass<"barrier-removal-continuation"> { +def SCFBarrierRemovalContinuation : Pass<"barrier-removal-continuation", "FuncOp"> { let summary = "Remove scf.barrier using continuations"; let constructor = "mlir::polygeist::createBarrierRemovalContinuation()"; let dependentDialects = ["memref::MemRefDialect", "StandardOpsDialect"]; } -def SCFRaiseToAffine : FunctionPass<"raise-scf-to-affine"> { +def SCFRaiseToAffine : Pass<"raise-scf-to-affine", "FuncOp"> { let summary = "Raise SCF to affine"; let constructor = "mlir::polygeist::createRaiseSCFToAffinePass()"; let dependentDialects = ["AffineDialect"]; } -def SCFCanonicalizeFor : FunctionPass<"canonicalize-scf-for"> { +def SCFCanonicalizeFor : Pass<"canonicalize-scf-for", "FuncOp"> { let summary = "Run some additional canonicalization for scf::for"; let constructor = "mlir::polygeist::createCanonicalizeForPass()"; } -def LoopRestructure : FunctionPass<"loop-restructure"> { +def LoopRestructure : Pass<"loop-restructure", "FuncOp"> { let constructor = "mlir::polygeist::createLoopRestructurePass()"; let dependentDialects = ["::mlir::scf::SCFDialect"]; } -def RemoveTrivialUse : FunctionPass<"trivialuse"> { +def RemoveTrivialUse : Pass<"trivialuse", "FuncOp"> { let constructor = "mlir::polygeist::createRemoveTrivialUsePass()"; } diff --git a/lib/polygeist/Ops.cpp b/lib/polygeist/Ops.cpp index 3608b186378c..9853cede9779 100644 --- a/lib/polygeist/Ops.cpp +++ b/lib/polygeist/Ops.cpp @@ -7,23 +7,23 @@ //===----------------------------------------------------------------------===// #include "polygeist/Ops.h" +#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/LLVMIR/LLVMTypes.h" #include "mlir/IR/Builders.h" #include "mlir/IR/OpImplementation.h" #include "mlir/Interfaces/SideEffectInterfaces.h" #include "polygeist/Dialect.h" -#include #define GET_OP_CLASSES #include "polygeist/PolygeistOps.cpp.inc" #include "mlir/Dialect/Affine/IR/AffineOps.h" +#include "mlir/Dialect/Arithmetic/Utils/Utils.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" +#include "mlir/Dialect/SCF/SCF.h" #include "mlir/Dialect/StandardOps/IR/Ops.h" -#include "mlir/Dialect/StandardOps/Utils/Utils.h" #include "mlir/IR/BlockAndValueMapping.h" -#include using namespace mlir; using namespace polygeist; @@ -161,8 +161,8 @@ class SubToCast final : public OpRewritePattern { if (cidx.value() != 0) return failure(); - rewriter.replaceOpWithNewOp(subViewOp, subViewOp.source(), - post); + rewriter.replaceOpWithNewOp(subViewOp, post, + subViewOp.source()); return success(); } @@ -724,7 +724,7 @@ MutableOperandRange LoadSelect::ptrMutable(LLVM::LoadOp op) { return op.getAddrMutable(); } -void SubIndexOp::getCanonicalizationPatterns(OwningRewritePatternList &results, +void SubIndexOp::getCanonicalizationPatterns(RewritePatternSet &results, MLIRContext *context) { results.insert { op.getLoc(), c0, rewriter.create( op.getLoc(), - rewriter.create(op.getLoc(), op.getLen(), - rewriter.getIndexType()), + rewriter.create( + op.getLoc(), rewriter.getIndexType(), op.getLen()), rewriter.create(op.getLoc(), width)), c1); @@ -918,8 +918,8 @@ OpFoldResult Memref2PointerOp::fold(ArrayRef operands) { return nullptr; } -void Memref2PointerOp::getCanonicalizationPatterns( - OwningRewritePatternList &results, MLIRContext *context) { +void Memref2PointerOp::getCanonicalizationPatterns(RewritePatternSet &results, + MLIRContext *context) { results.insert, CopySimplification>(context); @@ -1021,16 +1021,16 @@ class MetaPointer2Memref final : public OpRewritePattern { auto shape = mt.getShape(); for (size_t i = 0; i < shape.size(); i++) { auto off = computeIndex(op, i, rewriter); - auto cur = - rewriter.create(op.getLoc(), rewriter.getI32Type(), off); + auto cur = rewriter.create( + op.getLoc(), rewriter.getI32Type(), off); if (idx == nullptr) { idx = cur; } else { idx = rewriter.create( op.getLoc(), - rewriter.create( - op.getLoc(), idx, - rewriter.create(op.getLoc(), shape[i], 32)), + rewriter.create(op.getLoc(), idx, + rewriter.create( + op.getLoc(), shape[i], 32)), cur); } } @@ -1433,8 +1433,8 @@ struct MoveIntoIfs : public OpRewritePattern { } }; -void Pointer2MemrefOp::getCanonicalizationPatterns( - OwningRewritePatternList &results, MLIRContext *context) { +void Pointer2MemrefOp::getCanonicalizationPatterns(RewritePatternSet &results, + MLIRContext *context) { results.insert< Pointer2MemrefCast, Pointer2Memref2PointerCast, MetaPointer2Memref, MetaPointer2Memref, @@ -1517,7 +1517,7 @@ struct TypeSizeCanonicalize : public OpRewritePattern { } }; -void TypeSizeOp::getCanonicalizationPatterns(OwningRewritePatternList &results, +void TypeSizeOp::getCanonicalizationPatterns(RewritePatternSet &results, MLIRContext *context) { results.insert(context); } diff --git a/lib/polygeist/Passes/AffineCFG.cpp b/lib/polygeist/Passes/AffineCFG.cpp index f534e0aba467..72b1f8fc28ea 100644 --- a/lib/polygeist/Passes/AffineCFG.cpp +++ b/lib/polygeist/Passes/AffineCFG.cpp @@ -501,7 +501,7 @@ void fully2ComposeIntegerSetAndOperands(IntegerSet *set, namespace { struct AffineCFGPass : public AffineCFGBase { - void runOnFunction() override; + void runOnOperation() override; }; } // namespace @@ -561,9 +561,9 @@ struct SimplfyIntegerCastMath : public OpRewritePattern { setLocationAfter(b2, iadd.getOperand(1)); rewriter.replaceOpWithNewOp( op, - b.create(op.getLoc(), iadd.getOperand(0), op.getType()), - b2.create(op.getLoc(), iadd.getOperand(1), - op.getType())); + b.create(op.getLoc(), op.getType(), iadd.getOperand(0)), + b2.create(op.getLoc(), op.getType(), + iadd.getOperand(1))); return success(); } if (auto iadd = op.getOperand().getDefiningOp()) { @@ -573,9 +573,10 @@ struct SimplfyIntegerCastMath : public OpRewritePattern { setLocationAfter(b2, iadd.getOperand(1)); rewriter.replaceOpWithNewOp( op, - b.create(op.getLoc(), iadd.getOperand(0), op.getType()), - b2.create(op.getLoc(), iadd.getOperand(1), - op.getType())); + b.create(op.getLoc(), op.getType(), + iadd.getOperand(0)), + b2.create(op.getLoc(), op.getType(), + iadd.getOperand(1))); return success(); } if (auto iadd = op.getOperand().getDefiningOp()) { @@ -585,9 +586,9 @@ struct SimplfyIntegerCastMath : public OpRewritePattern { setLocationAfter(b2, iadd.getOperand(1)); rewriter.replaceOpWithNewOp( op, - b.create(op.getLoc(), iadd.getOperand(0), op.getType()), - b2.create(op.getLoc(), iadd.getOperand(1), - op.getType())); + b.create(op.getLoc(), op.getType(), iadd.getOperand(0)), + b2.create(op.getLoc(), op.getType(), + iadd.getOperand(1))); return success(); } if (auto iadd = op.getOperand().getDefiningOp()) { @@ -597,9 +598,10 @@ struct SimplfyIntegerCastMath : public OpRewritePattern { setLocationAfter(b2, iadd.getOperand(1)); rewriter.replaceOpWithNewOp( op, - b.create(op.getLoc(), iadd.getOperand(0), op.getType()), - b2.create(op.getLoc(), iadd.getOperand(1), - op.getType())); + b.create(op.getLoc(), op.getType(), + iadd.getOperand(0)), + b2.create(op.getLoc(), op.getType(), + iadd.getOperand(1))); return success(); } if (auto iadd = op.getOperand().getDefiningOp()) { @@ -609,9 +611,10 @@ struct SimplfyIntegerCastMath : public OpRewritePattern { setLocationAfter(b2, iadd.getOperand(1)); rewriter.replaceOpWithNewOp( op, - b.create(op.getLoc(), iadd.getOperand(0), op.getType()), - b2.create(op.getLoc(), iadd.getOperand(1), - op.getType())); + b.create(op.getLoc(), op.getType(), + iadd.getOperand(0)), + b2.create(op.getLoc(), op.getType(), + iadd.getOperand(1))); return success(); } return failure(); @@ -753,8 +756,8 @@ bool handle(OpBuilder &b, CmpIOp cmpi, SmallVectorImpl &exprs, } SmallVector lhspack = {cmpi.getLhs()}; if (!lhspack[0].getType().isa()) { - auto op = b.create(cmpi.getLoc(), lhspack[0], - IndexType::get(cmpi.getContext())); + auto op = b.create( + cmpi.getLoc(), IndexType::get(cmpi.getContext()), lhspack[0]); lhspack[0] = op; } @@ -762,8 +765,8 @@ bool handle(OpBuilder &b, CmpIOp cmpi, SmallVectorImpl &exprs, AffineMap::get(0, 1, getAffineSymbolExpr(0, cmpi.getContext())); SmallVector rhspack = {cmpi.getRhs()}; if (!rhspack[0].getType().isa()) { - auto op = b.create(cmpi.getLoc(), rhspack[0], - IndexType::get(cmpi.getContext())); + auto op = b.create( + cmpi.getLoc(), IndexType::get(cmpi.getContext()), rhspack[0]); rhspack[0] = op; } @@ -1131,16 +1134,15 @@ struct MoveIfToAffine : public OpRewritePattern { } }; -void AffineCFGPass::runOnFunction() { - mlir::RewritePatternSet rpl(getFunction().getContext()); +void AffineCFGPass::runOnOperation() { + mlir::RewritePatternSet rpl(getOperation().getContext()); rpl.add, AffineFixup, CanonicalizIfBounds, MoveStoreToAffine, MoveIfToAffine, MoveLoadToAffine, CanonicalieForBounds>( - getFunction().getContext()); + getOperation().getContext()); GreedyRewriteConfig config; - (void)applyPatternsAndFoldGreedily(getFunction().getOperation(), - std::move(rpl), config); + (void)applyPatternsAndFoldGreedily(getOperation(), std::move(rpl), config); } std::unique_ptr> mlir::polygeist::replaceAffineCFGPass() { diff --git a/lib/polygeist/Passes/AffineReduction.cpp b/lib/polygeist/Passes/AffineReduction.cpp index ce4862c130b2..304360c31ee9 100644 --- a/lib/polygeist/Passes/AffineReduction.cpp +++ b/lib/polygeist/Passes/AffineReduction.cpp @@ -13,7 +13,7 @@ using namespace polygeist; namespace { struct AffineReductionPass : public AffineReductionBase { - void runOnFunction() override; + void runOnOperation() override; }; } // end namespace. @@ -254,12 +254,11 @@ struct AffineForReductionIter : public OpRewritePattern { } // end namespace. -void AffineReductionPass::runOnFunction() { - mlir::RewritePatternSet rpl(getFunction().getContext()); - rpl.add(getFunction().getContext()); +void AffineReductionPass::runOnOperation() { + mlir::RewritePatternSet rpl(getOperation().getContext()); + rpl.add(getOperation().getContext()); GreedyRewriteConfig config; - (void)applyPatternsAndFoldGreedily(getFunction().getOperation(), - std::move(rpl), config); + (void)applyPatternsAndFoldGreedily(getOperation(), std::move(rpl), config); } namespace mlir { diff --git a/lib/polygeist/Passes/BarrierRemovalContinuation.cpp b/lib/polygeist/Passes/BarrierRemovalContinuation.cpp index c345107d57b9..f96001d19814 100644 --- a/lib/polygeist/Passes/BarrierRemovalContinuation.cpp +++ b/lib/polygeist/Passes/BarrierRemovalContinuation.cpp @@ -13,7 +13,9 @@ #include "PassDetails.h" -#include "mlir/Conversion/SCFToStandard/SCFToStandard.h" +#include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h" +#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/SCF/Passes.h" #include "mlir/Dialect/SCF/SCF.h" @@ -27,7 +29,6 @@ #include "mlir/Transforms/DialectConversion.h" #include "polygeist/BarrierUtils.h" #include "polygeist/Passes/Passes.h" -#include using namespace mlir; using namespace mlir::arith; @@ -72,7 +73,7 @@ static void wrapPersistingLoopBodies(FuncOp function) { /// Convert SCF constructs except parallel ops with immediate barriers to a CFG. static LogicalResult applyCFGConversion(FuncOp function) { RewritePatternSet patterns(function.getContext()); - populateLoopToStdConversionPatterns(patterns); + populateSCFToControlFlowConversionPatterns(patterns); // Configure the target to preserve parallel ops with barriers, unless those // barriers are nested in deeper parallel ops. @@ -104,7 +105,7 @@ static void splitBlocksWithBarrier(Region ®ion) { Block *original = op->getBlock(); Block *block = original->splitBlock(op->getNextNode()); auto builder = OpBuilder::atBlockEnd(original); - builder.create(builder.getUnknownLoc(), block); + builder.create(builder.getUnknownLoc(), block); } } @@ -225,8 +226,8 @@ replicateIntoRegion(Region ®ion, Value storage, ValueRange ivs, // Branch from the entry block to the first cloned block. builder.setInsertionPointToEnd(entryBlock); - builder.create(builder.getUnknownLoc(), - mapping.lookup(blocks.front())); + builder.create(builder.getUnknownLoc(), + mapping.lookup(blocks.front())); // Now that the block structure is created, clone the operations and introduce // the flow between continuations. @@ -244,7 +245,7 @@ replicateIntoRegion(Region ®ion, Value storage, ValueRange ivs, // blocks are assumed to branch to the entry block of another subgraph. // They are replaced with storing the correspnding continuation ID and a // yield. - if (auto branch = dyn_cast(&op)) { + if (auto branch = dyn_cast(&op)) { // if (!blocks.contains(branch.dest())) { if (isa_and_nonnull(branch->getPrevNode())) { auto it = llvm::find(subgraphEntryPoints, branch.getDest()); @@ -603,8 +604,8 @@ static void createContinuations(FuncOp func) { namespace { struct BarrierRemoval : public SCFBarrierRemovalContinuationBase { - void runOnFunction() override { - auto f = getFunction(); + void runOnOperation() override { + FuncOp f = getOperation(); if (failed(convertToCFG(f))) return; if (failed(splitBlocksWithBarrier(f))) diff --git a/lib/polygeist/Passes/CMakeLists.txt b/lib/polygeist/Passes/CMakeLists.txt index 19f5ec443855..9c10f356bf8b 100644 --- a/lib/polygeist/Passes/CMakeLists.txt +++ b/lib/polygeist/Passes/CMakeLists.txt @@ -31,6 +31,6 @@ add_mlir_dialect_library(MLIRPolygeistTransforms MLIRSideEffectInterfaces MLIRStandard MLIRStandardOpsTransforms - MLIRSCFToStandard + MLIRSCFToControlFlow MLIRTransformUtils ) diff --git a/lib/polygeist/Passes/CanonicalizeFor.cpp b/lib/polygeist/Passes/CanonicalizeFor.cpp index 2866cd0e581e..9f371af75575 100644 --- a/lib/polygeist/Passes/CanonicalizeFor.cpp +++ b/lib/polygeist/Passes/CanonicalizeFor.cpp @@ -17,7 +17,7 @@ using namespace polygeist; namespace { struct CanonicalizeFor : public SCFCanonicalizeForBase { - void runOnFunction() override; + void runOnOperation() override; }; } // namespace @@ -103,7 +103,7 @@ struct ForOpInductionReplacement : public OpRewritePattern { if (!std::get<1>(it).getType().isa()) { replacement = rewriter.create( - forOp.getLoc(), replacement, std::get<1>(it).getType()); + forOp.getLoc(), std::get<1>(it).getType(), replacement); } if (!sameValue) @@ -128,7 +128,7 @@ struct ForOpInductionReplacement : public OpRewritePattern { if (!std::get<1>(it).getType().isa()) { replacement = rewriter.create( - forOp.getLoc(), replacement, std::get<1>(it).getType()); + forOp.getLoc(), std::get<1>(it).getType(), replacement); } if (!sameValue) @@ -530,12 +530,12 @@ struct MoveWhileToFor : public OpRewritePattern { } } - Value ub = rewriter.create(loop.getLoc(), loopInfo.ub, - IndexType::get(loop.getContext())); - Value lb = rewriter.create(loop.getLoc(), loopInfo.lb, - IndexType::get(loop.getContext())); - step = rewriter.create(loop.getLoc(), step, - IndexType::get(loop.getContext())); + Value ub = rewriter.create( + loop.getLoc(), IndexType::get(loop.getContext()), loopInfo.ub); + Value lb = rewriter.create( + loop.getLoc(), IndexType::get(loop.getContext()), loopInfo.lb); + step = rewriter.create( + loop.getLoc(), IndexType::get(loop.getContext()), step); // input of the for goes the input of the scf::while plus the output taken // from the conditionOp. @@ -556,7 +556,7 @@ struct MoveWhileToFor : public OpRewritePattern { } else res = arg; if (cst) { - res = rewriter.create(rewriter.getUnknownLoc(), res, cst); + res = rewriter.create(rewriter.getUnknownLoc(), cst, res); } forArgs.push_back(res); } @@ -673,7 +673,8 @@ struct MoveWhileDown : public OpRewritePattern { rewriter.updateRootInPlace(op, [&] { for (auto val : todo) { - auto na = op.getAfter().front().addArgument(val.getType()); + auto na = + op.getAfter().front().addArgument(val.getType(), op->getLoc()); val.replaceUsesWithIf(na, [&](OpOperand &u) -> bool { return op.getAfter().isAncestor(u.getOwner()->getParentRegion()); }); @@ -853,7 +854,7 @@ struct MoveWhileDown2 : public OpRewritePattern { for (auto v : sv) { condArgs.push_back(v); - auto arg = afterB->addArgument(v.getType()); + auto arg = afterB->addArgument(v.getType(), ifOp->getLoc()); for (OpOperand &use : llvm::make_early_inc_range(v.getUses())) { if (ifOp->isAncestor(use.getOwner()) || use.getOwner() == afterYield) @@ -1124,7 +1125,8 @@ struct MoveWhileDown3 : public OpRewritePattern { llvm::make_early_inc_range(cloned->getOpOperands())) { { newOps.push_back(o.get()); - o.set(op.getAfter().front().addArgument(o.get().getType())); + o.set(op.getAfter().front().addArgument(o.get().getType(), + o.get().getLoc())); } } continue; @@ -1338,8 +1340,8 @@ struct MoveSideEffectFreeWhile : public OpRewritePattern { for (auto arg : term.getArgs()) { if (auto IC = arg.getDefiningOp()) { if (arg.hasOneUse() && op.getResult(i).use_empty()) { - auto rep = - op.getAfter().front().addArgument(IC->getOperand(0).getType()); + auto rep = op.getAfter().front().addArgument( + IC->getOperand(0).getType(), IC->getOperand(0).getLoc()); IC->moveBefore(&op.getAfter().front(), op.getAfter().front().begin()); conds.push_back(IC.getIn()); IC.getInMutable().assign(rep); @@ -1405,8 +1407,8 @@ struct ReturnSq : public OpRewritePattern { return success(changed); } }; -void CanonicalizeFor::runOnFunction() { - mlir::RewritePatternSet rpl(getFunction().getContext()); +void CanonicalizeFor::runOnOperation() { + mlir::RewritePatternSet rpl(getOperation().getContext()); rpl.add(getFunction().getContext()); + MoveSideEffectFreeWhile>(getOperation().getContext()); GreedyRewriteConfig config; config.maxIterations = 47; - (void)applyPatternsAndFoldGreedily(getFunction().getOperation(), - std::move(rpl), config); + (void)applyPatternsAndFoldGreedily(getOperation(), std::move(rpl), config); } std::unique_ptr mlir::polygeist::createCanonicalizeForPass() { diff --git a/lib/polygeist/Passes/ConvertPolygeistToLLVM.cpp b/lib/polygeist/Passes/ConvertPolygeistToLLVM.cpp index 0e2171e684f4..468824f9b231 100644 --- a/lib/polygeist/Passes/ConvertPolygeistToLLVM.cpp +++ b/lib/polygeist/Passes/ConvertPolygeistToLLVM.cpp @@ -13,6 +13,7 @@ #include "mlir/Analysis/DataLayoutAnalysis.h" #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h" +#include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h" #include "mlir/Conversion/LLVMCommon/ConversionTarget.h" #include "mlir/Conversion/LLVMCommon/Pattern.h" #include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h" @@ -24,7 +25,6 @@ #include "mlir/Dialect/StandardOps/IR/Ops.h" #include "mlir/Dialect/StandardOps/Transforms/Passes.h" #include "polygeist/Ops.h" - #define DEBUG_TYPE "convert-polygeist-to-llvm" using namespace mlir; @@ -349,11 +349,12 @@ struct ConvertPolygeistToLLVMPass LLVMTypeConverter converter(&getContext(), options, &dataLayoutAnalysis); RewritePatternSet patterns(&getContext()); populatePolygeistToLLVMConversionPatterns(converter, patterns); + cf::populateControlFlowToLLVMConversionPatterns(converter, patterns); populateMemRefToLLVMConversionPatterns(converter, patterns); populateStdToLLVMConversionPatterns(converter, patterns); populateOpenMPToLLVMConversionPatterns(converter, patterns); arith::populateArithmeticToLLVMConversionPatterns(converter, patterns); - populateStdExpandOpsPatterns(patterns); + patterns .add( converter); diff --git a/lib/polygeist/Passes/LoopRestructure.cpp b/lib/polygeist/Passes/LoopRestructure.cpp index 055db287d332..60d72c7addb2 100644 --- a/lib/polygeist/Passes/LoopRestructure.cpp +++ b/lib/polygeist/Passes/LoopRestructure.cpp @@ -9,6 +9,7 @@ //===----------------------------------------------------------------------===// #include "PassDetails.h" +#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/SCF/SCF.h" #include "mlir/Dialect/StandardOps/IR/Ops.h" @@ -188,7 +189,7 @@ struct LoopRestructure : public LoopRestructureBase { void runOnRegion(DominanceInfo &domInfo, Region ®ion); bool removeIfFromRegion(DominanceInfo &domInfo, Region ®ion, Block *pseudoExit); - void runOnFunction() override; + void runOnOperation() override; }; } // end anonymous namespace @@ -220,7 +221,7 @@ class LoopInfo : public llvm::LoopInfoBase { template class llvm::LoopBase; template class llvm::LoopInfoBase; -void LoopRestructure::runOnFunction() { +void LoopRestructure::runOnOperation() { // FuncOp f = getFunction(); DominanceInfo &domInfo = getAnalysis(); if (auto region = getOperation().getCallableRegion()) { @@ -231,7 +232,7 @@ void LoopRestructure::runOnFunction() { bool attemptToFoldIntoPredecessor(Block *target) { SmallVector P(target->pred_begin(), target->pred_end()); if (P.size() == 1) { - if (auto op = dyn_cast(P[0]->getTerminator())) { + if (auto op = dyn_cast(P[0]->getTerminator())) { assert(target->getNumArguments() == op.getNumOperands()); for (size_t i = 0; i < target->getNumArguments(); ++i) { target->getArgument(i).replaceAllUsesWith(op.getOperand(i)); @@ -243,7 +244,7 @@ bool attemptToFoldIntoPredecessor(Block *target) { return true; } } else if (P.size() == 2) { - if (auto op = dyn_cast(P[0]->getTerminator())) { + if (auto op = dyn_cast(P[0]->getTerminator())) { assert(target->getNumArguments() == op.getNumTrueOperands()); assert(target->getNumArguments() == op.getNumFalseOperands()); @@ -254,7 +255,7 @@ bool attemptToFoldIntoPredecessor(Block *target) { } for (size_t i = 0; i < target->getNumArguments(); ++i) { - auto sel = builder.create( + auto sel = builder.create( op.getLoc(), op.getCondition(), op.getTrueOperand(i), op.getFalseOperand(i)); target->getArgument(i).replaceAllUsesWith(sel); @@ -290,7 +291,7 @@ bool LoopRestructure::removeIfFromRegion(DominanceInfo &domInfo, Region ®ion, for (size_t j = 0; j < Succs.size(); ++j) { if (Succs[j] == pseudoExit && Succs[1 - j] == Preds[1 - i]) { OpBuilder builder(Preds[i]->getTerminator()); - auto condBr = cast(Preds[i]->getTerminator()); + auto condBr = cast(Preds[i]->getTerminator()); auto ifOp = builder.create( builder.getUnknownLoc(), condTys, condBr.getCondition(), /*hasElse*/ true); @@ -384,7 +385,9 @@ void LoopRestructure::runOnRegion(DominanceInfo &domInfo, Region ®ion) { // Copy the arguments across SmallVector headerArgumentTypes(header->getArgumentTypes()); - wrapper->addArguments(headerArgumentTypes); + SmallVector locs(headerArgumentTypes.size(), + builder.getUnknownLoc()); + wrapper->addArguments(headerArgumentTypes, locs); SmallVector valsCallingLoop; for (auto a : wrapper->getArguments()) @@ -404,7 +407,7 @@ void LoopRestructure::runOnRegion(DominanceInfo &domInfo, Region ®ion) { headerArgumentTypes.push_back(V.getType()); valsCallingLoop.push_back(builder.create( builder.getUnknownLoc(), V.getType())); - header->addArgument(V.getType()); + header->addArgument(V.getType(), V.getLoc()); } } } @@ -421,7 +424,7 @@ void LoopRestructure::runOnRegion(DominanceInfo &domInfo, Region ®ion) { for (size_t i = 0; i < returns.size(); ++i) { RetVals.push_back(loop.getResult(i + headerArgumentTypes.size())); } - builder.create(builder.getUnknownLoc(), target, RetVals); + builder.create(builder.getUnknownLoc(), target, RetVals); } for (auto &pair : preservedVals) { pair.first.replaceUsesWithIf(loop.getResult(pair.second), @@ -474,7 +477,8 @@ void LoopRestructure::runOnRegion(DominanceInfo &domInfo, Region ®ion) { Block *pseudoExit = new Block(); { insertRegion.push_back(pseudoExit); - pseudoExit->addArguments(tys); + SmallVector locs(tys.size(), builder.getUnknownLoc()); + pseudoExit->addArguments(tys, locs); OpBuilder builder(pseudoExit, pseudoExit->begin()); tys.clear(); builder.create(builder.getUnknownLoc(), tys, @@ -498,13 +502,13 @@ void LoopRestructure::runOnRegion(DominanceInfo &domInfo, Region ®ion) { for (auto v : preservedVals) args[v.second + 1] = v.first; - if (auto op = dyn_cast(terminator)) { + if (auto op = dyn_cast(terminator)) { args.insert(args.end(), op.getOperands().begin(), op.getOperands().end()); - builder.create(op.getLoc(), pseudoExit, args); + builder.create(op.getLoc(), pseudoExit, args); op.erase(); } - if (auto op = dyn_cast(terminator)) { + if (auto op = dyn_cast(terminator)) { std::vector trueargs(op.getTrueOperands().begin(), op.getTrueOperands().end()); std::vector falseargs(op.getFalseOperands().begin(), @@ -515,7 +519,7 @@ void LoopRestructure::runOnRegion(DominanceInfo &domInfo, Region ®ion) { if (op.getFalseDest() == target) { falseargs.insert(falseargs.begin(), args.begin(), args.end()); } - builder.create( + builder.create( op.getLoc(), op.getCondition(), op.getTrueDest() == target ? pseudoExit : op.getTrueDest(), trueargs, @@ -546,7 +550,7 @@ void LoopRestructure::runOnRegion(DominanceInfo &domInfo, Region ®ion) { auto vtrue = builder.create( builder.getUnknownLoc(), true, 1); - if (auto op = dyn_cast(terminator)) { + if (auto op = dyn_cast(terminator)) { SmallVector args(op.getOperands()); args.insert(args.begin(), vtrue); for (auto p : preservedVals) @@ -556,9 +560,9 @@ void LoopRestructure::runOnRegion(DominanceInfo &domInfo, Region ®ion) { builder.getUnknownLoc(), ty)); } terminator = - builder.create(op.getLoc(), pseudoExit, args); + builder.create(op.getLoc(), pseudoExit, args); op.erase(); - } else if (auto op = dyn_cast(terminator)) { + } else if (auto op = dyn_cast(terminator)) { std::vector trueargs(op.getTrueOperands().begin(), op.getTrueOperands().end()); std::vector falseargs(op.getFalseOperands().begin(), @@ -583,7 +587,7 @@ void LoopRestructure::runOnRegion(DominanceInfo &domInfo, Region ®ion) { } // Recreate the terminator and store it so that its other // successor is visited on the next iteration of the loop. - terminator = builder.create( + terminator = builder.create( op.getLoc(), op.getCondition(), op.getTrueDest() == header ? pseudoExit : op.getTrueDest(), trueargs, @@ -596,7 +600,8 @@ void LoopRestructure::runOnRegion(DominanceInfo &domInfo, Region ®ion) { } Block *after = new Block(); - after->addArguments(combinedTypes); + SmallVector locs2(combinedTypes.size(), region.getLoc()); + after->addArguments(combinedTypes, locs2); loop.getAfter().push_back(after); OpBuilder builder2(after, after->begin()); SmallVector yieldargs; @@ -623,9 +628,11 @@ void LoopRestructure::runOnRegion(DominanceInfo &domInfo, Region ®ion) { }); } - for (auto pair : - llvm::zip(header->getArguments(), - loopEntry->addArguments(header->getArgumentTypes()))) { + SmallVector locs3(header->getArgumentTypes().size(), + region.getLoc()); + for (auto pair : llvm::zip( + header->getArguments(), + loopEntry->addArguments(header->getArgumentTypes(), locs3))) { std::get<0>(pair).replaceAllUsesWith(std::get<1>(pair)); } header->eraseArguments([](BlockArgument) { return true; }); diff --git a/lib/polygeist/Passes/Mem2Reg.cpp b/lib/polygeist/Passes/Mem2Reg.cpp index 95fb64bc46cb..fdf0fca1e811 100644 --- a/lib/polygeist/Passes/Mem2Reg.cpp +++ b/lib/polygeist/Passes/Mem2Reg.cpp @@ -13,9 +13,10 @@ // SSA scalars live out of 'affine.for'/'affine.if' statements is available. //===----------------------------------------------------------------------===// #include "PassDetails.h" -#include "mlir/Analysis/AffineAnalysis.h" -#include "mlir/Analysis/Utils.h" +#include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" +#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/SCF/SCF.h" @@ -29,7 +30,6 @@ #include #include #include -#include #include #include "polygeist/Ops.h" @@ -76,7 +76,7 @@ namespace { // than dealloc) remain. // struct Mem2Reg : public Mem2RegBase { - void runOnFunction() override; + void runOnOperation() override; // return if changed bool forwardStoreToLoad(mlir::Value AI, std::vector idx, @@ -730,7 +730,7 @@ void removeRedundantBlockArgs( for (auto pred : prepred) { mlir::Value pval = nullptr; - if (auto op = dyn_cast(pred->getTerminator())) { + if (auto op = dyn_cast(pred->getTerminator())) { pval = op.getOperands()[blockArg.getArgNumber()]; if (pval.getType() != elType) { pval.getDefiningOp()->getParentRegion()->getParentOp()->dump(); @@ -739,7 +739,7 @@ void removeRedundantBlockArgs( assert(pval.getType() == elType); if (pval == blockArg) pval = nullptr; - } else if (auto op = dyn_cast(pred->getTerminator())) { + } else if (auto op = dyn_cast(pred->getTerminator())) { if (op.getTrueDest() == block) { if (blockArg.getArgNumber() >= op.getTrueOperands().size()) { block->dump(); @@ -766,7 +766,7 @@ void removeRedundantBlockArgs( if (pval == blockArg) pval = nullptr; } - } else if (auto op = dyn_cast(pred->getTerminator())) { + } else if (auto op = dyn_cast(pred->getTerminator())) { mlir::OpBuilder subbuilder(op.getOperation()); if (op.getDefaultDestination() == block) { pval = op.getDefaultOperands()[blockArg.getArgNumber()]; @@ -814,7 +814,7 @@ void removeRedundantBlockArgs( bool used = false; for (auto U : blockArg.getUsers()) { - if (auto op = dyn_cast(U)) { + if (auto op = dyn_cast(U)) { size_t i = 0; for (auto V : op.getOperands()) { if (V == blockArg && @@ -825,7 +825,7 @@ void removeRedundantBlockArgs( } if (used) break; - } else if (auto op = dyn_cast(U)) { + } else if (auto op = dyn_cast(U)) { size_t i = 0; for (auto V : op.getTrueOperands()) { if (V == blockArg && @@ -873,15 +873,16 @@ void removeRedundantBlockArgs( SetVector prepred(block->getPredecessors().begin(), block->getPredecessors().end()); for (auto pred : prepred) { - if (auto op = dyn_cast(pred->getTerminator())) { + if (auto op = dyn_cast(pred->getTerminator())) { mlir::OpBuilder subbuilder(op.getOperation()); std::vector args(op.getOperands().begin(), op.getOperands().end()); args.erase(args.begin() + blockArg.getArgNumber()); assert(args.size() == op.getOperands().size() - 1); - subbuilder.create(op.getLoc(), op.getDest(), args); + subbuilder.create(op.getLoc(), op.getDest(), args); op.erase(); - } else if (auto op = dyn_cast(pred->getTerminator())) { + } else if (auto op = + dyn_cast(pred->getTerminator())) { mlir::OpBuilder subbuilder(op.getOperation()); std::vector trueargs(op.getTrueOperands().begin(), @@ -896,11 +897,11 @@ void removeRedundantBlockArgs( } assert(trueargs.size() < op.getTrueOperands().size() || falseargs.size() < op.getFalseOperands().size()); - subbuilder.create(op.getLoc(), op.getCondition(), - op.getTrueDest(), trueargs, - op.getFalseDest(), falseargs); + subbuilder.create(op.getLoc(), op.getCondition(), + op.getTrueDest(), trueargs, + op.getFalseDest(), falseargs); op.erase(); - } else if (auto op = dyn_cast(pred->getTerminator())) { + } else if (auto op = dyn_cast(pred->getTerminator())) { mlir::OpBuilder builder(op.getOperation()); SmallVector defaultOps(op.getDefaultOperands().begin(), op.getDefaultOperands().end()); @@ -917,7 +918,7 @@ void removeRedundantBlockArgs( } vrange.push_back(cases.back()); } - builder.create( + builder.create( op.getLoc(), op.getFlag(), op.getDefaultDestination(), defaultOps, op.getCaseValuesAttr(), op.getCaseDestinations(), vrange); op.erase(); @@ -943,6 +944,7 @@ bool Mem2Reg::forwardStoreToLoad(mlir::Value AI, std::vector idx, bool changed = false; std::set loadOps; mlir::Type subType = nullptr; + mlir::Location loc = AI.getLoc(); std::set allStoreOps; std::deque> list = {{AI, false}}; @@ -1442,7 +1444,8 @@ bool Mem2Reg::forwardStoreToLoad(mlir::Value AI, std::vector idx, assert(endFind != valueAtEndOfBlock.end()); // Only handle known termination blocks - if (!isa(Pred->getTerminator())) { + if (!isa( + Pred->getTerminator())) { PotentialArgs[block] = Legality::Illegal; break; } @@ -1525,7 +1528,7 @@ bool Mem2Reg::forwardStoreToLoad(mlir::Value AI, std::vector idx, assert(startFound != valueAtStartOfBlock.end()); assert(startFound->second->valueAtStart == block); - auto arg = block->addArgument(subType); + auto arg = block->addArgument(subType, loc); auto argVal = metaMap.get(arg); valueAtStartOfBlock[block] = argVal; blocksWithAddedArgs[block] = arg; @@ -1578,14 +1581,14 @@ bool Mem2Reg::forwardStoreToLoad(mlir::Value AI, std::vector idx, assert(pred->getTerminator()); assert(blockArg.getOwner() == block); - if (auto op = dyn_cast(pred->getTerminator())) { + if (auto op = dyn_cast(pred->getTerminator())) { mlir::OpBuilder subbuilder(op.getOperation()); std::vector args(op.getOperands().begin(), op.getOperands().end()); args.push_back(pval); - subbuilder.create(op.getLoc(), op.getDest(), args); + subbuilder.create(op.getLoc(), op.getDest(), args); op.erase(); - } else if (auto op = dyn_cast(pred->getTerminator())) { + } else if (auto op = dyn_cast(pred->getTerminator())) { mlir::OpBuilder subbuilder(op.getOperation()); std::vector trueargs(op.getTrueOperands().begin(), @@ -1598,11 +1601,11 @@ bool Mem2Reg::forwardStoreToLoad(mlir::Value AI, std::vector idx, if (op.getFalseDest() == block) { falseargs.push_back(pval); } - subbuilder.create(op.getLoc(), op.getCondition(), - op.getTrueDest(), trueargs, - op.getFalseDest(), falseargs); + subbuilder.create(op.getLoc(), op.getCondition(), + op.getTrueDest(), trueargs, + op.getFalseDest(), falseargs); op.erase(); - } else if (auto op = dyn_cast(pred->getTerminator())) { + } else if (auto op = dyn_cast(pred->getTerminator())) { mlir::OpBuilder builder(op.getOperation()); SmallVector defaultOps(op.getDefaultOperands().begin(), op.getDefaultOperands().end()); @@ -1620,7 +1623,7 @@ bool Mem2Reg::forwardStoreToLoad(mlir::Value AI, std::vector idx, } vrange.push_back(cases.back()); } - builder.create( + builder.create( op.getLoc(), op.getFlag(), op.getDefaultDestination(), defaultOps, op.getCaseValuesAttr(), op.getCaseDestinations(), vrange); op.erase(); @@ -1779,9 +1782,9 @@ StoreMap getLastStored(mlir::Value AI) { return lastStored; } -void Mem2Reg::runOnFunction() { +void Mem2Reg::runOnOperation() { // Only supports single block functions at the moment. - FuncOp f = getFunction(); + FuncOp f = getOperation(); // Variable indicating that a memref has had a load removed // and or been deleted. Because there can be memrefs of diff --git a/lib/polygeist/Passes/ParallelLoopDistribute.cpp b/lib/polygeist/Passes/ParallelLoopDistribute.cpp index 832d722cf84a..5067336a0a61 100644 --- a/lib/polygeist/Passes/ParallelLoopDistribute.cpp +++ b/lib/polygeist/Passes/ParallelLoopDistribute.cpp @@ -197,12 +197,6 @@ static bool hasNestedBarrier(Operation *op, SmallVector &vals) { namespace { -/// Returns `true` if `value` is defined outside of the region that contains -/// `user`. -static bool isDefinedAbove(Value value, Operation *user) { - return value.getParentRegion()->isProperAncestor(user->getParentRegion()); -} - #if 0 /// Returns `true` if the loop has a form expected by interchange patterns. static bool isNormalized(scf::ForOp op) { @@ -541,7 +535,7 @@ static void moveBodiesIf(PatternRewriter &rewriter, T op, IfType ifOp, auto newParallel = rewriter.cloneWithoutRegions(op); newParallel.getRegion().push_back(new Block()); for (auto a : op.getBody()->getArguments()) - newParallel.getBody()->addArgument(a.getType()); + newParallel.getBody()->addArgument(a.getType(), op->getLoc()); rewriter.setInsertionPointToEnd(newParallel.getBody()); rewriter.clone(*op.getBody()->getTerminator()); @@ -565,7 +559,7 @@ static void moveBodiesIf(PatternRewriter &rewriter, T op, IfType ifOp, auto newParallel = rewriter.cloneWithoutRegions(op); newParallel.getRegion().push_back(new Block()); for (auto a : op.getBody()->getArguments()) - newParallel.getBody()->addArgument(a.getType()); + newParallel.getBody()->addArgument(a.getType(), op->getLoc()); rewriter.setInsertionPointToEnd(newParallel.getBody()); rewriter.clone(*op.getBody()->getTerminator()); @@ -695,7 +689,7 @@ static void moveBodiesFor(PatternRewriter &rewriter, T op, ForType forLoop, auto newParallel = rewriter.cloneWithoutRegions(op); newParallel.getRegion().push_back(new Block()); for (auto a : op.getBody()->getArguments()) - newParallel.getBody()->addArgument(a.getType()); + newParallel.getBody()->addArgument(a.getType(), op->getLoc()); rewriter.setInsertionPointToEnd(newParallel.getBody()); rewriter.clone(*op.getBody()->getTerminator()); @@ -890,11 +884,11 @@ template struct InterchangeWhilePFor : public OpRewritePattern { auto beforeParallelOp = rewriter.cloneWithoutRegions(op); beforeParallelOp.getRegion().push_back(new Block()); for (auto a : op.getBody()->getArguments()) - beforeParallelOp.getBody()->addArgument(a.getType()); + beforeParallelOp.getBody()->addArgument(a.getType(), a.getLoc()); auto afterParallelOp = rewriter.cloneWithoutRegions(op); afterParallelOp.getRegion().push_back(new Block()); for (auto a : op.getBody()->getArguments()) - afterParallelOp.getBody()->addArgument(a.getType()); + afterParallelOp.getBody()->addArgument(a.getType(), a.getLoc()); rewriter.setInsertionPointToEnd(beforeParallelOp.getBody()); rewriter.clone(*op.getBody()->getTerminator()); rewriter.setInsertionPointToEnd(afterParallelOp.getBody()); @@ -1167,7 +1161,7 @@ struct DistributeAroundBarrier : public OpRewritePattern { } idx = rewriter.create(ao.getLoc(), sz, rewriter.create( - ao.getLoc(), idx, sz.getType())); + ao.getLoc(), sz.getType(), idx)); SmallVector vec = {idx}; u.set(rewriter.create(ao.getLoc(), ao.getType(), alloc, idx)); @@ -1593,9 +1587,9 @@ struct Reg2MemWhile : public OpRewritePattern { struct CPUifyPass : public SCFCPUifyBase { CPUifyPass() = default; CPUifyPass(StringRef method) { this->method.setValue(method.str()); } - void runOnFunction() override { + void runOnOperation() override { if (method == "distribute") { - OwningRewritePatternList patterns(&getContext()); + RewritePatternSet patterns(&getContext()); patterns.insert, Reg2MemFor, Reg2MemWhile, Reg2MemIf, Reg2MemIf, WrapForWithBarrier, WrapAffineForWithBarrier, @@ -1628,12 +1622,12 @@ struct CPUifyPass : public SCFCPUifyBase { DistributeAroundBarrier>(&getContext()); GreedyRewriteConfig config; config.maxIterations = 142; - if (failed(applyPatternsAndFoldGreedily(getFunction(), + if (failed(applyPatternsAndFoldGreedily(getOperation(), std::move(patterns), config))) signalPassFailure(); } else if (method == "omp") { SmallVector toReplace; - getFunction().walk( + getOperation().walk( [&](polygeist::BarrierOp b) { toReplace.push_back(b); }); for (auto b : toReplace) { OpBuilder Builder(b); diff --git a/lib/polygeist/Passes/ParallelLower.cpp b/lib/polygeist/Passes/ParallelLower.cpp index a4578861d39a..d236450ae2c4 100644 --- a/lib/polygeist/Passes/ParallelLower.cpp +++ b/lib/polygeist/Passes/ParallelLower.cpp @@ -11,10 +11,10 @@ //===----------------------------------------------------------------------===// #include "PassDetails.h" -#include "mlir/Analysis/AffineAnalysis.h" #include "mlir/Analysis/CallGraph.h" -#include "mlir/Analysis/Utils.h" +#include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" +#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/GPU/GPUDialect.h" #include "mlir/Dialect/LLVMIR/NVVMDialect.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" @@ -125,7 +125,7 @@ struct AlwaysInlinerInterface : public InlinerInterface { // Replace the return with a branch to the dest. OpBuilder builder(op); - builder.create(op->getLoc(), newDest, returnOp.getOperands()); + builder.create(op->getLoc(), newDest, returnOp.getOperands()); op->erase(); } @@ -290,7 +290,7 @@ void ParallelLower::runOnOperation() { // mlir::OpBuilder builder2(f.getContext()); // builder2.setInsertionPointToStart(threadB); // iter++; - // builder2.create(loc, &*iter); + // builder2.create(loc, &*iter); container.walk([&](mlir::gpu::BlockIdOp bidx) { mlir::OpBuilder bz(launchOp.getContext()); @@ -320,7 +320,7 @@ void ParallelLower::runOnOperation() { alop.getType().getElementType(), alop.getType().getLayout(), Attribute())); alop.replaceAllUsesWith((mlir::Value)bz.create( - alop.getLoc(), newAlloca, alop.getType())); + alop.getLoc(), alop.getType(), newAlloca)); alop.erase(); } }); diff --git a/lib/polygeist/Passes/RaiseToAffine.cpp b/lib/polygeist/Passes/RaiseToAffine.cpp index 0479c211602a..3a52c84d8084 100644 --- a/lib/polygeist/Passes/RaiseToAffine.cpp +++ b/lib/polygeist/Passes/RaiseToAffine.cpp @@ -1,16 +1,16 @@ #include "PassDetails.h" #include "mlir/Dialect/Affine/IR/AffineOps.h" +#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/Dialect/SCF/Passes.h" #include "mlir/Dialect/SCF/SCF.h" +#include "mlir/Dialect/StandardOps/IR/Ops.h" #include "mlir/IR/BlockAndValueMapping.h" #include "mlir/Transforms/DialectConversion.h" #include "polygeist/Passes/Passes.h" #include "llvm/Support/Debug.h" -#include - #define DEBUG_TYPE "raise-to-affine" using namespace mlir; @@ -19,7 +19,7 @@ using namespace polygeist; namespace { struct RaiseSCFToAffine : public SCFRaiseToAffineBase { - void runOnFunction() override; + void runOnOperation() override; }; } // namespace @@ -219,16 +219,16 @@ struct ParallelOpRaising : public OpRewritePattern { } }; -void RaiseSCFToAffine::runOnFunction() { +void RaiseSCFToAffine::runOnOperation() { ConversionTarget target(getContext()); target .addLegalDialect(); - OwningRewritePatternList patterns(&getContext()); + RewritePatternSet patterns(&getContext()); patterns.insert(&getContext()); if (failed( - applyPartialConversion(getFunction(), target, std::move(patterns)))) + applyPartialConversion(getOperation(), target, std::move(patterns)))) signalPassFailure(); } diff --git a/lib/polygeist/Passes/TrivialUse.cpp b/lib/polygeist/Passes/TrivialUse.cpp index ac9b73e6202f..ffbbb6d464e6 100644 --- a/lib/polygeist/Passes/TrivialUse.cpp +++ b/lib/polygeist/Passes/TrivialUse.cpp @@ -21,7 +21,7 @@ using namespace polygeist; namespace { struct RemoveTrivialUse : public RemoveTrivialUseBase { - void runOnFunction() override; + void runOnOperation() override; }; } // end anonymous namespace @@ -34,6 +34,6 @@ std::unique_ptr> createRemoveTrivialUsePass() { } // namespace polygeist } // namespace mlir -void RemoveTrivialUse::runOnFunction() { +void RemoveTrivialUse::runOnOperation() { getOperation()->walk([&](polygeist::TrivialUseOp bidx) { bidx.erase(); }); } diff --git a/llvm-project b/llvm-project index b634f8a663d5..af6b9939aac0 160000 --- a/llvm-project +++ b/llvm-project @@ -1 +1 @@ -Subproject commit b634f8a663d56877663f5224a785d9f0263c4176 +Subproject commit af6b9939aac065ce19bdaaf9fc54a8368367f33f diff --git a/test/polygeist-opt/canonicalizefor.mlir b/test/polygeist-opt/canonicalizefor.mlir index bea14f6de817..d588306e9b5a 100644 --- a/test/polygeist-opt/canonicalizefor.mlir +++ b/test/polygeist-opt/canonicalizefor.mlir @@ -13,7 +13,7 @@ module { %5 = call @cmp() : () -> i1 scf.condition(%5) %arg3, %3, %2 : i32, index, index } do { - ^bb0(%arg3: i32, %arg4: index, %arg5: index): // no predecessors + ^bb0(%arg3: i32, %arg4: index, %arg5: index): %parg3 = arith.addi %arg3, %c1_i32 : i32 %3 = memref.load %arg0[%arg5] : memref memref.store %3, %arg1[%arg4] : memref @@ -33,7 +33,7 @@ module { // CHECK-NEXT: %1 = call @cmp() : () -> i1 // CHECK-NEXT: scf.condition(%1) %arg3 : i32 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg3: i32): // no predecessors +// CHECK-NEXT: ^bb0(%arg3: i32): // CHECK-NEXT: %1 = arith.index_cast %arg3 : i32 to index // CHECK-NEXT: %2 = arith.index_cast %arg3 : i32 to index // CHECK-NEXT: %3 = arith.addi %1, %c3 : index @@ -60,7 +60,7 @@ module { } scf.condition(%1) %2#0, %2#1 : i32, i32 } do { - ^bb0(%arg2: i32, %arg3: i32): // no predecessors + ^bb0(%arg2: i32, %arg3: i32): scf.yield %arg2, %arg3 : i32, i32 } return %0#1 : i32 @@ -73,7 +73,7 @@ module { // CHECK-NEXT: %1 = arith.cmpi sgt, %arg2, %c0_i32 : i32 // CHECK-NEXT: scf.condition(%1) %arg3, %arg2 : i32, i32 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg2: i32, %arg3: i32): // no predecessors +// CHECK-NEXT: ^bb0(%arg2: i32, %arg3: i32): // CHECK-NEXT: %1 = arith.remsi %arg2, %arg3 : i32 // CHECK-NEXT: scf.yield %1, %arg3 : i32, i32 // CHECK-NEXT: } @@ -98,7 +98,7 @@ module { } scf.condition(%1) %2 : i32 } do { - ^bb0(%arg2: i32): // no predecessors + ^bb0(%arg2: i32): scf.yield %arg2 : i32 } return %c0_i32 : i32 @@ -136,7 +136,7 @@ module { } scf.condition(%2) %3#0, %3#1, %3#2 : f32, i32, i32 } do { - ^bb0(%arg2: f32, %arg3: i32, %arg4: i32): // no predecessors + ^bb0(%arg2: f32, %arg3: i32, %arg4: i32): scf.yield %arg2, %arg3, %arg4 : f32, i32, i32 } return %0#1 : i32 @@ -151,7 +151,7 @@ module { // CHECK-NEXT: %1 = arith.cmpf ult, %arg2, %arg0 : f32 // CHECK-NEXT: scf.condition(%1) %arg3, %arg2, %arg4 : i32, f32, i32 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg2: i32, %arg3: f32, %arg4: i32): // no predecessors +// CHECK-NEXT: ^bb0(%arg2: i32, %arg3: f32, %arg4: i32): // CHECK-NEXT: %1 = arith.addf %arg3, %arg1 : f32 // CHECK-NEXT: scf.yield %1, %arg4, %arg2 : f32, i32, i32 // CHECK-NEXT: } diff --git a/test/polygeist-opt/cudalower.mlir b/test/polygeist-opt/cudalower.mlir index e44c1d18f515..e012c6421218 100644 --- a/test/polygeist-opt/cudalower.mlir +++ b/test/polygeist-opt/cudalower.mlir @@ -26,13 +26,13 @@ module attributes {llvm.data_layout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64 module { func private @S(%arg0: i8, %arg1: !llvm.ptr) -> i8 { - switch %arg0 : i8, [ + cf.switch %arg0 : i8, [ default: ^bb10(%arg0 : i8), 0: ^bb1 ] ^bb1: // 2 preds: ^bb0, ^bb0 %6 = llvm.load %arg1 : !llvm.ptr - br ^bb10(%6 : i8) + cf.br ^bb10(%6 : i8) ^bb10(%50: i8): // 10 preds: ^bb0, ^bb1, ^bb2, ^bb3, ^bb4, ^bb5, ^bb6, ^bb7, ^bb8, ^bb9 return %50 : i8 } @@ -54,15 +54,15 @@ module { // CHECK-NEXT: scf.parallel (%arg2, %arg3, %arg4) = (%c0, %c0, %c0) to (%c2, %c1, %c1) step (%c1, %c1, %c1) { // CHECK-NEXT: scf.parallel (%arg5, %arg6, %arg7) = (%c0, %c0, %c0) to (%c1, %c1, %c1) step (%c1, %c1, %c1) { // CHECK-NEXT: %0 = scf.execute_region -> i8 { -// CHECK-NEXT: switch %arg1 : i8, [ +// CHECK-NEXT: cf.switch %arg1 : i8, [ // CHECK-NEXT: default: ^bb2(%arg1 : i8), // CHECK-NEXT: 0: ^bb1 // CHECK-NEXT: ] // CHECK-NEXT: ^bb1: // pred: ^bb0 // CHECK-NEXT: %1 = llvm.load %arg0 : !llvm.ptr -// CHECK-NEXT: br ^bb2(%1 : i8) +// CHECK-NEXT: cf.br ^bb2(%1 : i8) // CHECK-NEXT: ^bb2(%2: i8): // 2 preds: ^bb0, ^bb1 -// CHECK-NEXT: br ^bb3(%2 : i8) +// CHECK-NEXT: cf.br ^bb3(%2 : i8) // CHECK-NEXT: ^bb3(%3: i8): // pred: ^bb2 // CHECK-NEXT: scf.yield %3 : i8 // CHECK-NEXT: } diff --git a/test/polygeist-opt/fath.mlir b/test/polygeist-opt/fath.mlir index 1af7ccf58adf..7a9411b33bc8 100644 --- a/test/polygeist-opt/fath.mlir +++ b/test/polygeist-opt/fath.mlir @@ -4,13 +4,13 @@ module { func @_Z14computeTempCPUPfS_iii(%arg2: i32, %arg3: i32, %arg4: i32) { %c1_i32 = arith.constant 1 : i32 %c0_i32 = arith.constant 0 : i32 - br ^bb1(%c0_i32, %c0_i32 : i32, i32) + cf.br ^bb1(%c0_i32, %c0_i32 : i32, i32) ^bb1(%0: i32, %1: i32): // 3 preds: ^bb0, ^bb4, ^bb5 %2 = arith.cmpi ult, %0, %arg3 : i32 - cond_br %2, ^bb2, ^bb5 + cf.cond_br %2, ^bb2, ^bb5 ^bb2: %9 = arith.addi %0, %c1_i32 : i32 - br ^bb1(%9, %1 : i32, i32) + cf.br ^bb1(%9, %1 : i32, i32) ^bb5: // pred: ^bb1 return } @@ -31,7 +31,7 @@ module { // CHECK-NEXT: } // CHECK-NEXT: scf.condition(%2#0) %2#1, %2#2 : i32, i32 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg3: i32, %arg4: i32): // no predecessors +// CHECK-NEXT: ^bb0(%arg3: i32, %arg4: i32): // CHECK-NEXT: scf.yield %arg3, %arg4 : i32, i32 // CHECK-NEXT: } // CHECK-NEXT: return diff --git a/test/polygeist-opt/fath2.mlir b/test/polygeist-opt/fath2.mlir index 2c7c5c2c8b79..073f804980f5 100644 --- a/test/polygeist-opt/fath2.mlir +++ b/test/polygeist-opt/fath2.mlir @@ -4,27 +4,27 @@ module { func @_Z14computeTempCPUPfS_iii(%arg0: memref, %arg1: memref, %arg2: i32, %arg3: i32, %arg4: i32) { %c1_i32 = arith.constant 1 : i32 %c0_i32 = arith.constant 0 : i32 - br ^bb1(%c0_i32, %c0_i32 : i32, i32) + cf.br ^bb1(%c0_i32, %c0_i32 : i32, i32) ^bb1(%0: i32, %1: i32): // 3 preds: ^bb0, ^bb4, ^bb5 %2 = arith.cmpi ult, %0, %arg3 : i32 - cond_br %2, ^bb2(%c0_i32 : i32), ^bb5 + cf.cond_br %2, ^bb2(%c0_i32 : i32), ^bb5 ^bb2(%3: i32): // 2 preds: ^bb1, ^bb3 %4 = arith.cmpi ult, %3, %arg2 : i32 - cond_br %4, ^bb3, ^bb4 + cf.cond_br %4, ^bb3, ^bb4 ^bb3: // pred: ^bb2 %5 = arith.index_cast %3 : i32 to index %6 = arith.index_cast %3 : i32 to index %7 = memref.load %arg0[%6] : memref memref.store %7, %arg1[%5] : memref %8 = arith.addi %3, %c1_i32 : i32 - br ^bb2(%8 : i32) + cf.br ^bb2(%8 : i32) ^bb4: // pred: ^bb2 %9 = arith.addi %0, %c1_i32 : i32 - br ^bb1(%9, %1 : i32, i32) + cf.br ^bb1(%9, %1 : i32, i32) ^bb5: // pred: ^bb1 %10 = arith.addi %1, %c1_i32 : i32 %11 = arith.cmpi ult, %10, %arg4 : i32 - cond_br %11, ^bb1(%c0_i32, %10 : i32, i32), ^bb6 + cf.cond_br %11, ^bb1(%c0_i32, %10 : i32, i32), ^bb6 ^bb6: // pred: ^bb5 return } @@ -36,13 +36,13 @@ module { // CHECK-NEXT: %0:2 = scf.while (%arg5 = %c0_i32, %arg6 = %c0_i32) : (i32, i32) -> (i32, i32) { // CHECK-NEXT: %1:3 = scf.execute_region -> (i1, i32, i32) { // CHECK-NEXT: %2 = arith.cmpi ult, %arg5, %arg3 : i32 -// CHECK-NEXT: cond_br %2, ^bb3(%c0_i32 : i32), ^bb1 +// CHECK-NEXT: cf.cond_br %2, ^bb3(%c0_i32 : i32), ^bb1 // CHECK-NEXT: ^bb1: // pred: ^bb0 // CHECK-NEXT: %3 = arith.addi %arg6, %c1_i32 : i32 // CHECK-NEXT: %4 = arith.cmpi ult, %3, %arg4 : i32 // CHECK-NEXT: %false = arith.constant false // CHECK-NEXT: %true = arith.constant true -// CHECK-NEXT: cond_br %4, ^bb2(%true, %c0_i32, %3 : i1, i32, i32), ^bb2(%false, %arg5, %arg6 : i1, i32, i32) +// CHECK-NEXT: cf.cond_br %4, ^bb2(%true, %c0_i32, %3 : i1, i32, i32), ^bb2(%false, %arg5, %arg6 : i1, i32, i32) // CHECK-NEXT: ^bb2(%5: i1, %6: i32, %7: i32): // 3 preds: ^bb1, ^bb1, ^bb3 // CHECK-NEXT: scf.yield %5, %6, %7 : i1, i32, i32 // CHECK-NEXT: ^bb3(%8: i32): // pred: ^bb0 @@ -62,16 +62,16 @@ module { // CHECK-NEXT: } // CHECK-NEXT: scf.condition(%12#0) %12#1 : i32 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg7: i32): // no predecessors +// CHECK-NEXT: ^bb0(%arg7: i32): // CHECK-NEXT: scf.yield %arg7 : i32 // CHECK-NEXT: } // CHECK-NEXT: %10 = arith.addi %arg5, %c1_i32 : i32 // CHECK-NEXT: %true_0 = arith.constant true -// CHECK-NEXT: br ^bb2(%true_0, %10, %arg6 : i1, i32, i32) +// CHECK-NEXT: cf.br ^bb2(%true_0, %10, %arg6 : i1, i32, i32) // CHECK-NEXT: } // CHECK-NEXT: scf.condition(%1#0) %1#1, %1#2 : i32, i32 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg5: i32, %arg6: i32): // no predecessors +// CHECK-NEXT: ^bb0(%arg5: i32, %arg6: i32): // CHECK-NEXT: scf.yield %arg5, %arg6 : i32, i32 // CHECK-NEXT: } // CHECK-NEXT: return diff --git a/test/polygeist-opt/induction.mlir b/test/polygeist-opt/induction.mlir index c77693b4f48f..a2d600548ac4 100644 --- a/test/polygeist-opt/induction.mlir +++ b/test/polygeist-opt/induction.mlir @@ -8,21 +8,21 @@ module { %c20 = arith.constant 20 : index %0 = memref.alloca() : memref memref.store %arg2, %0[] : memref - br ^bb02 + cf.br ^bb02 ^bb02: // pred: ^bb1 scf.if %arg1 { memref.store %c0, %0[] : memref } - br ^bb1 + cf.br ^bb1 ^bb1: // 2 preds: ^bb0, ^bb2 %7 = memref.load %0[] : memref %8 = arith.cmpi slt, %7, %c20 : index - cond_br %8, ^bb2, ^bb3 + cf.cond_br %8, ^bb2, ^bb3 ^bb2: // pred: ^bb1 memref.store %c0_i32, %arg0[%7] : memref<20xi32> %24 = arith.addi %7, %c1 : index memref.store %24, %0[] : memref - br ^bb1 + cf.br ^bb1 ^bb3: // pred: ^bb1 return } @@ -33,21 +33,21 @@ module { // CHECK-NEXT: %c0 = arith.constant 0 : index // CHECK-NEXT: %c1 = arith.constant 1 : index // CHECK-NEXT: %c20 = arith.constant 20 : index -// CHECK-NEXT: br ^bb1 +// CHECK-NEXT: cf.br ^bb1 // CHECK-NEXT: ^bb1: // pred: ^bb0 // CHECK-NEXT: %0 = scf.if %arg1 -> (index) { // CHECK-NEXT: scf.yield %c0 : index // CHECK-NEXT: } else { // CHECK-NEXT: scf.yield %arg2 : index // CHECK-NEXT: } -// CHECK-NEXT: br ^bb2(%0 : index) +// CHECK-NEXT: cf.br ^bb2(%0 : index) // CHECK-NEXT: ^bb2(%1: index): // 2 preds: ^bb1, ^bb3 // CHECK-NEXT: %2 = arith.cmpi slt, %1, %c20 : index -// CHECK-NEXT: cond_br %2, ^bb3, ^bb4 +// CHECK-NEXT: cf.cond_br %2, ^bb3, ^bb4 // CHECK-NEXT: ^bb3: // pred: ^bb2 // CHECK-NEXT: memref.store %c0_i32, %arg0[%1] : memref<20xi32> // CHECK-NEXT: %3 = arith.addi %1, %c1 : index -// CHECK-NEXT: br ^bb2(%3 : index) +// CHECK-NEXT: cf.br ^bb2(%3 : index) // CHECK-NEXT: ^bb4: // pred: ^bb2 // CHECK-NEXT: return // CHECK-NEXT: } diff --git a/test/polygeist-opt/infmem2ref.mlir b/test/polygeist-opt/infmem2ref.mlir index d40f9128eca6..bddd42289d71 100644 --- a/test/polygeist-opt/infmem2ref.mlir +++ b/test/polygeist-opt/infmem2ref.mlir @@ -13,12 +13,12 @@ module { %6 = memref.alloca() : memref memref.store %true, %6[] : memref scf.execute_region { - br ^bb1(%c0_i32 : i32) + cf.br ^bb1(%c0_i32 : i32) ^bb1(%28 : i32): // 2 preds: ^bb0, ^bb2 %29 = arith.cmpi slt, %28, %c2_i32 : i32 %30 = memref.load %6[] : memref %33 = arith.andi %29, %30 : i1 - cond_br %33, ^bb2, ^bb3 + cf.cond_br %33, ^bb2, ^bb3 ^bb2: // pred: ^bb1 scf.if %c { %45 = arith.cmpi eq, %28, %c1_i32 : i32 @@ -29,7 +29,7 @@ module { %44 = memref.load %6[] : memref call @use(%44) : (i1) -> () %42 = arith.addi %28, %c1_i32 : i32 - br ^bb1(%42 : i32) + cf.br ^bb1(%42 : i32) ^bb3: // pred: ^bb1 scf.yield } @@ -43,11 +43,11 @@ module { // CHECK-NEXT: %true = arith.constant true // CHECK-NEXT: %false = arith.constant false // CHECK-NEXT: scf.execute_region { -// CHECK-NEXT: br ^bb1(%c0_i32, %true : i32, i1) +// CHECK-NEXT: cf.br ^bb1(%c0_i32, %true : i32, i1) // CHECK-NEXT: ^bb1(%0: i32, %1: i1): // 2 preds: ^bb0, ^bb2 // CHECK-NEXT: %2 = arith.cmpi slt, %0, %c2_i32 : i32 // CHECK-NEXT: %3 = arith.andi %2, %1 : i1 -// CHECK-NEXT: cond_br %3, ^bb2, ^bb3 +// CHECK-NEXT: cf.cond_br %3, ^bb2, ^bb3 // CHECK-NEXT: ^bb2: // pred: ^bb1 // CHECK-NEXT: %4 = scf.if %arg0 -> (i1) { // CHECK-NEXT: %6 = arith.cmpi eq, %0, %c1_i32 : i32 @@ -62,7 +62,7 @@ module { // CHECK-NEXT: } // CHECK-NEXT: call @use(%4) : (i1) -> () // CHECK-NEXT: %5 = arith.addi %0, %c1_i32 : i32 -// CHECK-NEXT: br ^bb1(%5, %4 : i32, i1) +// CHECK-NEXT: cf.br ^bb1(%5, %4 : i32, i1) // CHECK-NEXT: ^bb3: // pred: ^bb1 // CHECK-NEXT: scf.yield // CHECK-NEXT: } @@ -78,12 +78,12 @@ module { %6 = memref.alloca() : memref memref.store %true, %6[] : memref scf.execute_region { - br ^bb1(%c0_i32 : i32) + cf.br ^bb1(%c0_i32 : i32) ^bb1(%28 : i32): // 2 preds: ^bb0, ^bb2 %29 = arith.cmpi slt, %28, %c2_i32 : i32 %30 = memref.load %6[] : memref %33 = arith.andi %29, %30 : i1 - cond_br %33, ^bb2, ^bb3 + cf.cond_br %33, ^bb2, ^bb3 ^bb2: // pred: ^bb1 scf.if %true { %45 = arith.cmpi eq, %28, %c1_i32 : i32 @@ -94,7 +94,7 @@ module { %44 = memref.load %6[] : memref call @use(%44) : (i1) -> () %42 = arith.addi %28, %c1_i32 : i32 - br ^bb1(%42 : i32) + cf.br ^bb1(%42 : i32) ^bb3: // pred: ^bb1 scf.yield } @@ -110,12 +110,12 @@ module { // CHECK-NEXT: %0 = memref.alloca() : memref // CHECK-NEXT: memref.store %true, %0[] : memref // CHECK-NEXT: scf.execute_region { -// CHECK-NEXT: br ^bb1(%c0_i32 : i32) +// CHECK-NEXT: cf.br ^bb1(%c0_i32 : i32) // CHECK-NEXT: ^bb1(%1: i32): // 2 preds: ^bb0, ^bb2 // CHECK-NEXT: %2 = arith.cmpi slt, %1, %c2_i32 : i32 // CHECK-NEXT: %3 = memref.load %0[] : memref // CHECK-NEXT: %4 = arith.andi %2, %3 : i1 -// CHECK-NEXT: cond_br %4, ^bb2, ^bb3 +// CHECK-NEXT: cf.cond_br %4, ^bb2, ^bb3 // CHECK-NEXT: ^bb2: // pred: ^bb1 // CHECK-NEXT: scf.if %true { // CHECK-NEXT: %7 = arith.cmpi eq, %1, %c1_i32 : i32 @@ -126,7 +126,7 @@ module { // CHECK-NEXT: %5 = memref.load %0[] : memref // CHECK-NEXT: call @use(%5) : (i1) -> () // CHECK-NEXT: %6 = arith.addi %1, %c1_i32 : i32 -// CHECK-NEXT: br ^bb1(%6 : i32) +// CHECK-NEXT: cf.br ^bb1(%6 : i32) // CHECK-NEXT: ^bb3: // pred: ^bb1 // CHECK-NEXT: scf.yield // CHECK-NEXT: } @@ -143,15 +143,15 @@ module { %1 = llvm.mlir.undef : i32 %2 = memref.alloca() : memref memref.store %1, %2[] : memref - br ^bb3 + cf.br ^bb3 ^bb3: // pred: ^bb1 scf.if %arg1 { memref.store %c1_i32, %2[] : memref scf.execute_region { - br ^bb1(%c0 : index) + cf.br ^bb1(%c0 : index) ^bb1(%11 : index): // 2 preds: ^bb0, ^bb2 %12 = arith.cmpi slt, %11, %c10 : index - cond_br %12, ^bb2, ^bb3 + cf.cond_br %12, ^bb2, ^bb3 ^bb2: // pred: ^bb1 %14 = memref.load %arg0[%11] : memref<11xf32> %15 = arith.cmpf ugt, %14, %cst : f32 @@ -159,7 +159,7 @@ module { memref.store %c0_i32, %2[] : memref } %16 = arith.addi %11, %c1 : index - br ^bb1(%16 : index) + cf.br ^bb1(%16 : index) ^bb3: // pred: ^bb1 scf.yield } diff --git a/test/polygeist-opt/mem2regRedundantArg.mlir b/test/polygeist-opt/mem2regRedundantArg.mlir index bdb502bf02da..82c577bf5b6e 100644 --- a/test/polygeist-opt/mem2regRedundantArg.mlir +++ b/test/polygeist-opt/mem2regRedundantArg.mlir @@ -9,16 +9,16 @@ module { %c10 = arith.constant 10 : index %2 = memref.alloca() : memref memref.store %true, %2[] : memref - br ^bb1(%c0 : index) + cf.br ^bb1(%c0 : index) ^bb1(%3 : index): // 2 preds: ^bb0, ^bb2 %4 = arith.cmpi slt, %3, %c10 : index %5 = memref.load %2[] : memref %6 = arith.andi %4, %5 : i1 - cond_br %6, ^bb2, ^bb3 + cf.cond_br %6, ^bb2, ^bb3 ^bb2: // pred: ^bb1 memref.store %cst, %arg0[%3] : memref %10 = arith.addi %3, %c1 : index - br ^bb1(%10 : index) + cf.br ^bb1(%10 : index) ^bb3: // pred: ^bb1 return } @@ -30,15 +30,15 @@ module { // CHECK-NEXT: %c0 = arith.constant 0 : index // CHECK-NEXT: %c1 = arith.constant 1 : index // CHECK-NEXT: %c10 = arith.constant 10 : index -// CHECK-NEXT: br ^bb1(%c0 : index) +// CHECK-NEXT: cf.br ^bb1(%c0 : index) // CHECK-NEXT: ^bb1(%0: index): // 2 preds: ^bb0, ^bb2 // CHECK-NEXT: %1 = arith.cmpi slt, %0, %c10 : index // CHECK-NEXT: %2 = arith.andi %1, %true : i1 -// CHECK-NEXT: cond_br %2, ^bb2, ^bb3 +// CHECK-NEXT: cf.cond_br %2, ^bb2, ^bb3 // CHECK-NEXT: ^bb2: // pred: ^bb1 // CHECK-NEXT: memref.store %cst, %arg0[%0] : memref // CHECK-NEXT: %3 = arith.addi %0, %c1 : index -// CHECK-NEXT: br ^bb1(%3 : index) +// CHECK-NEXT: cf.br ^bb1(%3 : index) // CHECK-NEXT: ^bb3: // pred: ^bb1 // CHECK-NEXT: return // CHECK-NEXT: } diff --git a/test/polygeist-opt/mem2regnest.mlir b/test/polygeist-opt/mem2regnest.mlir index a0ab0b8babdb..484ec8724203 100644 --- a/test/polygeist-opt/mem2regnest.mlir +++ b/test/polygeist-opt/mem2regnest.mlir @@ -9,7 +9,7 @@ module { %6 = llvm.mlir.undef : i32 memref.store %6, %5[] : memref %a9 = scf.execute_region -> (i32) { - cond_br %arg0, ^bb1, ^bb3 + cf.cond_br %arg0, ^bb1, ^bb3 ^bb1: // pred: ^bb0 scf.yield %c20_i32 : i32 ^bb3: // 2 preds: ^bb0, ^bb2 @@ -30,7 +30,7 @@ module { // CHECK-NEXT: %c10_i32 = arith.constant 10 : i32 // CHECK-NEXT: %0 = llvm.mlir.undef : i32 // CHECK-NEXT: %1:2 = scf.execute_region -> (i32, i32) { -// CHECK-NEXT: cond_br %arg0, ^bb1, ^bb2 +// CHECK-NEXT: cf.cond_br %arg0, ^bb1, ^bb2 // CHECK-NEXT: ^bb1: // pred: ^bb0 // CHECK-NEXT: scf.yield %c20_i32, %0 : i32, i32 // CHECK-NEXT: ^bb2: // pred: ^bb0 diff --git a/test/polygeist-opt/meme2regelse.mlir b/test/polygeist-opt/meme2regelse.mlir index def2dcf4f8af..9f25e1c986ee 100644 --- a/test/polygeist-opt/meme2regelse.mlir +++ b/test/polygeist-opt/meme2regelse.mlir @@ -6,12 +6,12 @@ module { %c0_i64 = arith.constant 0 : i64 %0 = memref.alloca() : memref memref.store %c0_i64, %0[] : memref - br ^bb1 + cf.br ^bb1 ^bb1: // 2 preds: ^bb0, ^bb2 scf.if %arg1 { } else { scf.execute_region { - cond_br %arg0, ^bb2, ^bb3 + cf.cond_br %arg0, ^bb2, ^bb3 ^bb2: // pred: ^bb1 memref.store %c1_i64, %0[] : memref scf.yield @@ -19,7 +19,7 @@ module { scf.yield } } - br ^bb3 + cf.br ^bb3 ^bb3: // pred: ^bb1 %8 = memref.load %0[] : memref return %8 : i64 @@ -29,13 +29,13 @@ module { // CHECK: func @bad(%arg0: i1, %arg1: i1, %arg2: memref) -> i64 { // CHECK-NEXT: %c1_i64 = arith.constant 1 : i64 // CHECK-NEXT: %c0_i64 = arith.constant 0 : i64 -// CHECK-NEXT: br ^bb1 +// CHECK-NEXT: cf.br ^bb1 // CHECK-NEXT: ^bb1: // pred: ^bb0 // CHECK-NEXT: %0 = scf.if %arg1 -> (i64) { // CHECK-NEXT: scf.yield %c0_i64 : i64 // CHECK-NEXT: } else { // CHECK-NEXT: %1 = scf.execute_region -> i64 { -// CHECK-NEXT: cond_br %arg0, ^bb1, ^bb2 +// CHECK-NEXT: cf.cond_br %arg0, ^bb1, ^bb2 // CHECK-NEXT: ^bb1: // pred: ^bb0 // CHECK-NEXT: scf.yield %c1_i64 : i64 // CHECK-NEXT: ^bb2: // pred: ^bb0 @@ -43,7 +43,7 @@ module { // CHECK-NEXT: } // CHECK-NEXT: scf.yield %1 : i64 // CHECK-NEXT: } -// CHECK-NEXT: br ^bb2 +// CHECK-NEXT: cf.br ^bb2 // CHECK-NEXT: ^bb2: // pred: ^bb1 // CHECK-NEXT: return %0 : i64 // CHECK-NEXT: } diff --git a/test/polygeist-opt/restructure.mlir b/test/polygeist-opt/restructure.mlir index 4aba7c8ba53a..f3f65271a220 100644 --- a/test/polygeist-opt/restructure.mlir +++ b/test/polygeist-opt/restructure.mlir @@ -4,14 +4,14 @@ module { func @kernel_gemm(%arg0: i64) -> i1 { %c0_i64 = arith.constant 0 : i64 %c1_i64 = arith.constant 1 : i64 - br ^bb1(%c0_i64 : i64) + cf.br ^bb1(%c0_i64 : i64) ^bb1(%0: i64): // 2 preds: ^bb0, ^bb2 %2 = arith.cmpi "slt", %0, %c0_i64 : i64 %5 = arith.cmpi "sle", %0, %arg0 : i64 - cond_br %5, ^bb2, ^bb3 + cf.cond_br %5, ^bb2, ^bb3 ^bb2: // pred: ^bb1 %8 = arith.addi %0, %c1_i64 : i64 - br ^bb1(%8 : i64) + cf.br ^bb1(%8 : i64) ^bb3: // pred: ^bb1 return %2 : i1 } @@ -34,7 +34,7 @@ func @kernel_gemm(%arg0: i64) -> i1 { // CHECK-NEXT: } // CHECK-NEXT: scf.condition(%4#0) %4#1, %4#2 : i64, i1 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg1: i64, %arg2: i1): // no predecessors +// CHECK-NEXT: ^bb0(%arg1: i64, %arg2: i1): // CHECK-NEXT: scf.yield %arg1, %arg2 : i64, i1 // CHECK-NEXT: } // CHECK-NEXT: return %1#1 : i1 @@ -49,11 +49,11 @@ func @kernel_gemm(%arg0: i64) -> i1 { %2 = memref.alloca() : memref memref.store %arg0, %2[] : memref memref.store %arg1, %1[] : memref - br ^bb1 + cf.br ^bb1 ^bb1: // 2 preds: ^bb0, ^bb2 %3 = memref.load %1[] : memref %4 = arith.cmpi sgt, %3, %c0_i32 : i32 - cond_br %4, ^bb2, ^bb3 + cf.cond_br %4, ^bb2, ^bb3 ^bb2: // pred: ^bb1 %5 = memref.load %0[] : memref %8 = memref.load %2[] : memref @@ -63,7 +63,7 @@ func @kernel_gemm(%arg0: i64) -> i1 { } memref.store %3, %2[] : memref memref.store %9, %1[] : memref - br ^bb1 + cf.br ^bb1 ^bb3: // pred: ^bb1 %7 = memref.load %2[] : memref return %7 : i32 diff --git a/tools/mlir-clang/CMakeLists.txt b/tools/mlir-clang/CMakeLists.txt index 7ec692b941a4..b54c1436f2f2 100644 --- a/tools/mlir-clang/CMakeLists.txt +++ b/tools/mlir-clang/CMakeLists.txt @@ -50,7 +50,7 @@ target_link_libraries(mlir-clang PRIVATE MLIROpenMP MLIRGPUOps MLIRTransforms - MLIRSCFToStandard + MLIRSCFToControlFlow MLIRStandardToLLVM MLIRAffineTransforms MLIRAffineToStandard diff --git a/tools/mlir-clang/Lib/CGCall.cc b/tools/mlir-clang/Lib/CGCall.cc index 56289eb3634a..ad90223ec7bc 100644 --- a/tools/mlir-clang/Lib/CGCall.cc +++ b/tools/mlir-clang/Lib/CGCall.cc @@ -41,8 +41,8 @@ static mlir::Value castCallerMemRefArg(mlir::Value callerArg, std::next(dstShape.begin()))) { b.setInsertionPointAfterValue(callerArg); - return b.create(callerArg.getLoc(), callerArg, - calleeArgType); + return b.create(callerArg.getLoc(), calleeArgType, + callerArg); } } } @@ -154,10 +154,11 @@ ValueCategory MLIRScanner::CallHelper( .store(builder, arg, /*isArray*/ isArray); shape[0] = pshape; val = builder.create( - loc, alloc, + loc, mlir::MemRefType::get(shape, mt.getElementType(), MemRefLayoutAttrInterface(), - mt.getMemorySpace())); + mt.getMemorySpace()), + alloc); } else { val = arg.getValue(builder); if (val.getType().isa() && @@ -168,7 +169,7 @@ ValueCategory MLIRScanner::CallHelper( if (auto prevTy = val.getType().dyn_cast()) { auto ipostTy = expectedType.cast(); if (prevTy != ipostTy) - val = builder.create(loc, val, ipostTy); + val = builder.create(loc, ipostTy, val); } } } else { @@ -235,10 +236,10 @@ ValueCategory MLIRScanner::CallHelper( mt.getMemorySpace())); shape[0] = pshape; alloc = builder.create( - loc, alloc, + loc, mlir::MemRefType::get(shape, mt.getElementType(), - MemRefLayoutAttrInterface(), - mt.getMemorySpace())); + MemRefLayoutAttrInterface(), mt.getMemorySpace()), + alloc); args.push_back(alloc); } @@ -252,21 +253,20 @@ ValueCategory MLIRScanner::CallHelper( mlir::Value idx[] = {getConstantIndex(0), getConstantIndex(i)}; assert(MT.getShape().size() == 2); blocks[i] = builder.create( - loc, builder.create(loc, val, idx), - mlir::IndexType::get(builder.getContext())); + loc, mlir::IndexType::get(builder.getContext()), + builder.create(loc, val, idx)); } else { mlir::Value idx[] = {builder.create(loc, 0, 32), builder.create(loc, i, 32)}; auto PT = val.getType().cast(); auto ET = PT.getElementType().cast().getBody()[i]; blocks[i] = builder.create( - loc, + loc, mlir::IndexType::get(builder.getContext()), builder.create( - loc, + loc, mlir::IndexType::get(builder.getContext()), builder.create( loc, LLVM::LLVMPointerType::get(ET, PT.getAddressSpace()), - val, idx)), - mlir::IndexType::get(builder.getContext())); + val, idx))); } } @@ -279,21 +279,20 @@ ValueCategory MLIRScanner::CallHelper( mlir::Value idx[] = {getConstantIndex(0), getConstantIndex(i)}; assert(MT.getShape().size() == 2); threads[i] = builder.create( - loc, builder.create(loc, val, idx), - mlir::IndexType::get(builder.getContext())); + loc, mlir::IndexType::get(builder.getContext()), + builder.create(loc, val, idx)); } else { mlir::Value idx[] = {builder.create(loc, 0, 32), builder.create(loc, i, 32)}; auto PT = val.getType().cast(); auto ET = PT.getElementType().cast().getBody()[i]; threads[i] = builder.create( - loc, + loc, mlir::IndexType::get(builder.getContext()), builder.create( loc, builder.create( loc, LLVM::LLVMPointerType::get(ET, PT.getAddressSpace()), - val, idx)), - mlir::IndexType::get(builder.getContext())); + val, idx))); } } auto op = builder.create(loc, blocks[0], blocks[1], @@ -671,7 +670,7 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) { : CmpFPredicate::ONE; mlir::Value FCmp = builder.create(loc, Pred, Fabs, Infinity); auto postTy = getMLIRType(expr->getType()).cast(); - mlir::Value res = builder.create(loc, FCmp, postTy); + mlir::Value res = builder.create(loc, postTy, FCmp); return ValueCategory(res, /*isRef*/ false); } if (sr->getDecl()->getIdentifier() && @@ -680,7 +679,7 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) { mlir::Value V = getLLVM(expr->getArg(0)); mlir::Value Eq = builder.create(loc, CmpFPredicate::UNO, V, V); auto postTy = getMLIRType(expr->getType()).cast(); - mlir::Value res = builder.create(loc, Eq, postTy); + mlir::Value res = builder.create(loc, postTy, Eq); return ValueCategory(res, /*isRef*/ false); } if (sr->getDecl()->getIdentifier() && @@ -702,7 +701,7 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) { V = builder.create(loc, Eq, IsLessThanInf); V = builder.create(loc, V, IsNormal); auto postTy = getMLIRType(expr->getType()).cast(); - mlir::Value res = builder.create(loc, V, postTy); + mlir::Value res = builder.create(loc, postTy, V); return ValueCategory(res, /*isRef*/ false); } if (sr->getDecl()->getIdentifier() && @@ -714,7 +713,7 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) { auto ZeroV = builder.create(loc, 0, ITy); V = builder.create(loc, CmpIPredicate::slt, BC, ZeroV); auto postTy = getMLIRType(expr->getType()).cast(); - mlir::Value res = builder.create(loc, V, postTy); + mlir::Value res = builder.create(loc, postTy, V); return ValueCategory(res, /*isRef*/ false); } if (sr->getDecl()->getIdentifier() && @@ -723,7 +722,7 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) { mlir::Value V2 = getLLVM(expr->getArg(1)); V = builder.create(loc, CmpFPredicate::OGT, V, V2); auto postTy = getMLIRType(expr->getType()).cast(); - mlir::Value res = builder.create(loc, V, postTy); + mlir::Value res = builder.create(loc, postTy, V); return ValueCategory(res, /*isRef*/ false); } if (sr->getDecl()->getIdentifier() && @@ -732,7 +731,7 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) { mlir::Value V2 = getLLVM(expr->getArg(1)); V = builder.create(loc, CmpFPredicate::OGE, V, V2); auto postTy = getMLIRType(expr->getType()).cast(); - mlir::Value res = builder.create(loc, V, postTy); + mlir::Value res = builder.create(loc, postTy, V); return ValueCategory(res, /*isRef*/ false); } if (sr->getDecl()->getIdentifier() && @@ -741,7 +740,7 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) { mlir::Value V2 = getLLVM(expr->getArg(1)); V = builder.create(loc, CmpFPredicate::OLT, V, V2); auto postTy = getMLIRType(expr->getType()).cast(); - mlir::Value res = builder.create(loc, V, postTy); + mlir::Value res = builder.create(loc, postTy, V); return ValueCategory(res, /*isRef*/ false); } if (sr->getDecl()->getIdentifier() && @@ -750,7 +749,7 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) { mlir::Value V2 = getLLVM(expr->getArg(1)); V = builder.create(loc, CmpFPredicate::OLE, V, V2); auto postTy = getMLIRType(expr->getType()).cast(); - mlir::Value res = builder.create(loc, V, postTy); + mlir::Value res = builder.create(loc, postTy, V); return ValueCategory(res, /*isRef*/ false); } if (sr->getDecl()->getIdentifier() && @@ -759,7 +758,7 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) { mlir::Value V2 = getLLVM(expr->getArg(1)); V = builder.create(loc, CmpFPredicate::ONE, V, V2); auto postTy = getMLIRType(expr->getType()).cast(); - mlir::Value res = builder.create(loc, V, postTy); + mlir::Value res = builder.create(loc, postTy, V); return ValueCategory(res, /*isRef*/ false); } if (sr->getDecl()->getIdentifier() && @@ -768,7 +767,7 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) { mlir::Value V2 = getLLVM(expr->getArg(1)); V = builder.create(loc, CmpFPredicate::UNO, V, V2); auto postTy = getMLIRType(expr->getType()).cast(); - mlir::Value res = builder.create(loc, V, postTy); + mlir::Value res = builder.create(loc, postTy, V); return ValueCategory(res, /*isRef*/ false); } if (sr->getDecl()->getIdentifier() && diff --git a/tools/mlir-clang/Lib/CGStmt.cc b/tools/mlir-clang/Lib/CGStmt.cc index 758f8c286020..8f96e5cf721d 100644 --- a/tools/mlir-clang/Lib/CGStmt.cc +++ b/tools/mlir-clang/Lib/CGStmt.cc @@ -9,6 +9,7 @@ #include "IfScope.h" #include "clang-mlir.h" #include "mlir/Dialect/Arithmetic/IR/Arithmetic.h" +#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h" #include "mlir/Dialect/OpenMP/OpenMPDialect.h" #include "mlir/Dialect/SCF/SCF.h" #include "mlir/IR/Diagnostics.h" @@ -54,7 +55,7 @@ bool MLIRScanner::getLowerBound(clang::ForStmt *fors, if (auto declRefStmt = dyn_cast(binOp->getLHS())) { mlir::Value val = Visit(binOp->getRHS()).getValue(builder); val = builder.create( - loc, val, mlir::IndexType::get(builder.getContext())); + loc, mlir::IndexType::get(builder.getContext()), val); descr.setName(cast(declRefStmt->getDecl())); descr.setType(getMLIRType(declRefStmt->getDecl()->getType())); if (descr.getForwardMode()) @@ -97,8 +98,8 @@ bool MLIRScanner::getUpperBound(clang::ForStmt *fors, auto rhs = binaryOp->getRHS(); mlir::Value val = Visit(rhs).getValue(builder); - val = builder.create(loc, val, - mlir::IndexType::get(val.getContext())); + val = builder.create( + loc, mlir::IndexType::get(val.getContext()), val); if (binaryOp->getOpcode() == clang::BinaryOperator::Opcode::BO_LE) val = builder.create(loc, val, getConstantIndex(1)); descr.setUpperBound(val); @@ -110,8 +111,8 @@ bool MLIRScanner::getUpperBound(clang::ForStmt *fors, auto rhs = binaryOp->getRHS(); mlir::Value val = Visit(rhs).getValue(builder); - val = builder.create(loc, val, - mlir::IndexType::get(val.getContext())); + val = builder.create( + loc, mlir::IndexType::get(val.getContext()), val); if (binaryOp->getOpcode() == clang::BinaryOperator::Opcode::BO_GT) val = builder.create(loc, val, getConstantIndex(1)); descr.setLowerBound(val); @@ -184,7 +185,7 @@ void MLIRScanner::buildAffineLoopImpl( val = builder.create( loc, builder.create(loc, ub, getConstantIndex(1)), val); } - auto idx = builder.create(loc, val, descr.getType()); + auto idx = builder.create(loc, descr.getType(), val); assert(params.find(descr.getName()) != params.end()); params[descr.getName()].store(builder, idx); @@ -238,7 +239,7 @@ ValueCategory MLIRScanner::VisitForStmt(clang::ForStmt *fors) { auto &exitB = *(new Block()); toadd->getBlocks().push_back(&exitB); - builder.create(loc, &condB); + builder.create(loc, &condB); builder.setInsertionPointToStart(&condB); @@ -259,7 +260,7 @@ ValueCategory MLIRScanner::VisitForStmt(clang::ForStmt *fors) { auto nb = builder.create( loc, lctx.noBreak, std::vector()); cond = builder.create(loc, cond, nb); - builder.create(loc, cond, &bodyB, &exitB); + builder.create(loc, cond, &bodyB, &exitB); } builder.setInsertionPointToStart(&bodyB); @@ -277,7 +278,7 @@ ValueCategory MLIRScanner::VisitForStmt(clang::ForStmt *fors) { loops.pop_back(); if (builder.getInsertionBlock()->empty() || !isTerminator(&builder.getInsertionBlock()->back())) { - builder.create(loc, &condB); + builder.create(loc, &condB); } builder.setInsertionPointToStart(&exitB); @@ -329,15 +330,15 @@ ValueCategory MLIRScanner::VisitOMPForDirective(clang::OMPForDirective *fors) { for (auto f : fors->inits()) { assert(f); f = cast(f)->getRHS(); - inits.push_back(builder.create(loc, Visit(f).getValue(builder), - builder.getIndexType())); + inits.push_back(builder.create(loc, builder.getIndexType(), + Visit(f).getValue(builder))); } SmallVector finals; for (auto f : fors->finals()) { f = cast(f)->getRHS(); - finals.push_back(builder.create( - loc, Visit(f).getValue(builder), builder.getIndexType())); + finals.push_back(builder.create(loc, builder.getIndexType(), + Visit(f).getValue(builder))); } SmallVector incs; @@ -353,14 +354,14 @@ ValueCategory MLIRScanner::VisitOMPForDirective(clang::OMPForDirective *fors) { bo = cast(f); assert(bo->getOpcode() == clang::BinaryOperator::Opcode::BO_Mul); f = bo->getRHS(); - incs.push_back(builder.create(loc, Visit(f).getValue(builder), - builder.getIndexType())); + incs.push_back(builder.create(loc, builder.getIndexType(), + Visit(f).getValue(builder))); } auto affineOp = builder.create(loc, inits, finals, incs); affineOp.getRegion().push_back(new Block()); for (auto init : inits) - affineOp.getRegion().front().addArgument(init.getType()); + affineOp.getRegion().front().addArgument(init.getType(), init.getLoc()); auto inds = affineOp.getRegion().front().getArguments(); auto oldpoint = builder.getInsertionPoint(); @@ -380,8 +381,8 @@ ValueCategory MLIRScanner::VisitOMPForDirective(clang::OMPForDirective *fors) { std::map prevInduction; for (auto zp : zip(inds, fors->counters())) { auto idx = builder.create( - loc, std::get<0>(zp), - getMLIRType(fors->getIterationVariable()->getType())); + loc, getMLIRType(fors->getIterationVariable()->getType()), + std::get<0>(zp)); VarDecl *name = cast(cast(std::get<1>(zp))->getDecl()); @@ -502,15 +503,15 @@ ValueCategory MLIRScanner::VisitOMPParallelForDirective( for (auto f : fors->inits()) { assert(f); f = cast(f)->getRHS(); - inits.push_back(builder.create(loc, Visit(f).getValue(builder), - builder.getIndexType())); + inits.push_back(builder.create(loc, builder.getIndexType(), + Visit(f).getValue(builder))); } SmallVector finals; for (auto f : fors->finals()) { f = cast(f)->getRHS(); - finals.push_back(builder.create( - loc, Visit(f).getValue(builder), builder.getIndexType())); + finals.push_back(builder.create( + loc, builder.getIndexType(), Visit(f).getValue(builder))); } SmallVector incs; @@ -526,8 +527,8 @@ ValueCategory MLIRScanner::VisitOMPParallelForDirective( bo = cast(f); assert(bo->getOpcode() == clang::BinaryOperator::Opcode::BO_Mul); f = bo->getRHS(); - incs.push_back(builder.create(loc, Visit(f).getValue(builder), - builder.getIndexType())); + incs.push_back(builder.create(loc, builder.getIndexType(), + Visit(f).getValue(builder))); } auto affineOp = builder.create(loc, inits, finals, incs); @@ -550,8 +551,8 @@ ValueCategory MLIRScanner::VisitOMPParallelForDirective( std::map prevInduction; for (auto zp : zip(inds, fors->counters())) { auto idx = builder.create( - loc, std::get<0>(zp), - getMLIRType(fors->getIterationVariable()->getType())); + loc, getMLIRType(fors->getIterationVariable()->getType()), + std::get<0>(zp)); VarDecl *name = cast(cast(std::get<1>(zp))->getDecl()); @@ -612,7 +613,7 @@ ValueCategory MLIRScanner::VisitDoStmt(clang::DoStmt *fors) { auto &exitB = *(new Block()); toadd->getBlocks().push_back(&exitB); - builder.create(loc, &bodyB); + builder.create(loc, &bodyB); builder.setInsertionPointToStart(&condB); @@ -633,7 +634,7 @@ ValueCategory MLIRScanner::VisitDoStmt(clang::DoStmt *fors) { auto nb = builder.create(loc, loops.back().noBreak, std::vector()); cond = builder.create(loc, cond, nb); - builder.create(loc, cond, &bodyB, &exitB); + builder.create(loc, cond, &bodyB, &exitB); } builder.setInsertionPointToStart(&bodyB); @@ -646,7 +647,7 @@ ValueCategory MLIRScanner::VisitDoStmt(clang::DoStmt *fors) { Visit(fors->getBody()); loops.pop_back(); - builder.create(loc, &condB); + builder.create(loc, &condB); builder.setInsertionPointToStart(&exitB); @@ -673,7 +674,7 @@ ValueCategory MLIRScanner::VisitWhileStmt(clang::WhileStmt *fors) { auto &exitB = *(new Block()); toadd->getBlocks().push_back(&exitB); - builder.create(loc, &condB); + builder.create(loc, &condB); builder.setInsertionPointToStart(&condB); @@ -694,7 +695,7 @@ ValueCategory MLIRScanner::VisitWhileStmt(clang::WhileStmt *fors) { auto nb = builder.create(loc, loops.back().noBreak, std::vector()); cond = builder.create(loc, cond, nb); - builder.create(loc, cond, &bodyB, &exitB); + builder.create(loc, cond, &bodyB, &exitB); } builder.setInsertionPointToStart(&bodyB); @@ -707,7 +708,7 @@ ValueCategory MLIRScanner::VisitWhileStmt(clang::WhileStmt *fors) { Visit(fors->getBody()); loops.pop_back(); - builder.create(loc, &condB); + builder.create(loc, &condB); builder.setInsertionPointToStart(&exitB); @@ -800,7 +801,7 @@ ValueCategory MLIRScanner::VisitSwitchStmt(clang::SwitchStmt *stmt) { if (inCase) { auto noBreak = builder.create(loc, loops.back().noBreak); - builder.create(loc, noBreak, &condB, &exitB); + builder.create(loc, noBreak, &condB, &exitB); loops.pop_back(); } @@ -824,7 +825,7 @@ ValueCategory MLIRScanner::VisitSwitchStmt(clang::SwitchStmt *stmt) { if (inCase) { auto noBreak = builder.create(loc, loops.back().noBreak); - builder.create(loc, noBreak, &condB, &exitB); + builder.create(loc, noBreak, &condB, &exitB); loops.pop_back(); } @@ -856,7 +857,7 @@ ValueCategory MLIRScanner::VisitSwitchStmt(clang::SwitchStmt *stmt) { if (inCase) loops.pop_back(); - builder.create(loc, &exitB); + builder.create(loc, &exitB); er.getRegion().getBlocks().push_back(&exitB); @@ -885,7 +886,7 @@ ValueCategory MLIRScanner::VisitSwitchStmt(clang::SwitchStmt *stmt) { } builder.setInsertionPointToStart(&er.getRegion().front()); - builder.create( + builder.create( loc, cond, defaultB, ArrayRef(), caseValuesAttr, blocks, SmallVector(caseVals.size(), ArrayRef())); builder.setInsertionPoint(oldblock2, oldpoint2); @@ -956,7 +957,7 @@ ValueCategory MLIRScanner::VisitLabelStmt(clang::LabelStmt *stmt) { labels[stmt] = labelB; } toadd->getBlocks().push_back(labelB); - builder.create(loc, labelB); + builder.create(loc, labelB); builder.setInsertionPointToStart(labelB); Visit(stmt->getSubStmt()); return nullptr; @@ -972,7 +973,7 @@ ValueCategory MLIRScanner::VisitGotoStmt(clang::GotoStmt *stmt) { labelB = new Block(); labels[labelstmt] = labelB; } - builder.create(loc, labelB); + builder.create(loc, labelB); return nullptr; } @@ -1022,7 +1023,7 @@ ValueCategory MLIRScanner::VisitReturnStmt(clang::ReturnStmt *stmt) { if (auto prevTy = val.getType().dyn_cast()) { auto ipostTy = postTy.cast(); if (prevTy != ipostTy) { - val = builder.create(loc, val, ipostTy); + val = builder.create(loc, ipostTy, val); } } else if (val.getType().isa() && postTy.isa()) diff --git a/tools/mlir-clang/Lib/TypeUtils.cc b/tools/mlir-clang/Lib/TypeUtils.cc index 402567ddd393..af9fe4dc48b4 100644 --- a/tools/mlir-clang/Lib/TypeUtils.cc +++ b/tools/mlir-clang/Lib/TypeUtils.cc @@ -33,7 +33,7 @@ bool isRecursiveStruct(Type *T, Type *Meta, SmallPtrSetImpl &seen) { Type *anonymize(Type *T) { if (auto PT = dyn_cast(T)) - return PointerType::get(anonymize(PT->getElementType()), + return PointerType::get(anonymize(PT->getPointerElementType()), PT->getAddressSpace()); if (auto AT = dyn_cast(T)) return ArrayType::get(anonymize(AT->getElementType()), diff --git a/tools/mlir-clang/Lib/clang-mlir.cc b/tools/mlir-clang/Lib/clang-mlir.cc index 7f721c5d7957..64c6735ccc4d 100644 --- a/tools/mlir-clang/Lib/clang-mlir.cc +++ b/tools/mlir-clang/Lib/clang-mlir.cc @@ -338,7 +338,7 @@ mlir::Value MLIRScanner::createAllocOp(mlir::Type t, VarDecl *name, varLoc, LLVM::LLVMPointerType::get(t, 0), alloc)); } alloc = abuilder.create( - varLoc, alloc, mlir::MemRefType::get(-1, t, {}, 0)); + varLoc, mlir::MemRefType::get(-1, t, {}, 0), alloc); if (t.isa()) { mlir::Value idxs[] = {abuilder.create(loc, 0)}; abuilder.create( @@ -359,7 +359,7 @@ mlir::Value MLIRScanner::createAllocOp(mlir::Type t, VarDecl *name, shape, mt.getElementType(), MemRefLayoutAttrInterface(), wrapIntegerMemorySpace(memspace, mt.getContext())); auto len = Visit(var->getSizeExpr()).getValue(builder); - len = builder.create(varLoc, len, builder.getIndexType()); + len = builder.create(varLoc, builder.getIndexType(), len); alloc = builder.create(varLoc, mr, len); builder.create(varLoc, alloc); if (memspace != 0) { @@ -387,7 +387,7 @@ mlir::Value MLIRScanner::createAllocOp(mlir::Type t, VarDecl *name, } shape[0] = pshape; alloc = abuilder.create( - varLoc, alloc, mlir::MemRefType::get(shape, mt.getElementType())); + varLoc, mlir::MemRefType::get(shape, mt.getElementType()), alloc); } } assert(alloc); @@ -1061,7 +1061,7 @@ ValueCategory MLIRScanner::VisitCXXNewExpr(clang::CXXNewExpr *expr) { if (expr->isArray()) { count = Visit(*expr->raw_arg_begin()).getValue(builder); count = builder.create( - loc, count, mlir::IndexType::get(builder.getContext())); + loc, mlir::IndexType::get(builder.getContext()), count); } else { count = getConstantIndex(1); } @@ -1079,7 +1079,7 @@ ValueCategory MLIRScanner::VisitCXXNewExpr(clang::CXXNewExpr *expr) { auto i64 = mlir::IntegerType::get(count.getContext(), 64); auto typeSize = getTypeSize(expr->getAllocatedType()); mlir::Value args[1] = {builder.create(loc, typeSize, count)}; - args[0] = builder.create(loc, args[0], i64); + args[0] = builder.create(loc, i64, args[0]); alloc = builder.create( loc, ty, builder @@ -1141,8 +1141,8 @@ mlir::Value MLIRScanner::castToIndex(mlir::Location loc, mlir::Value val) { if (auto op = val.getDefiningOp()) return getConstantIndex(op.value()); - return builder.create(loc, val, - mlir::IndexType::get(val.getContext())); + return builder.create( + loc, mlir::IndexType::get(val.getContext()), val); } ValueCategory @@ -1225,7 +1225,7 @@ ValueCategory MLIRScanner::VisitConstructCommon(clang::CXXConstructExpr *cons, auto i8_0 = builder.create(loc, 0, 8); auto sizev = - builder.create(loc, size, builder.getI64Type()); + builder.create(loc, builder.getI64Type(), size); auto falsev = builder.create(loc, false, 1); builder.create(loc, val, i8_0, sizev, falsev); @@ -1325,7 +1325,7 @@ ValueCategory MLIRScanner::CommonArrayLookup(ValueCategory array, if (val.getType().isa()) { mlir::Value vals[] = { - builder.create(loc, idx, builder.getIntegerType(64))}; + builder.create(loc, builder.getIntegerType(64), idx)}; // TODO sub return ValueCategory( builder.create(loc, val.getType(), val, vals), @@ -1590,7 +1590,7 @@ MLIRScanner::EmitGPUCallExpr(clang::CallExpr *expr) { allocSize = Visit(expr->getArg(1)).getValue(builder); auto idxType = mlir::IndexType::get(builder.getContext()); mlir::Value args[1] = {builder.create( - loc, builder.create(loc, allocSize, idxType), + loc, builder.create(loc, idxType, allocSize), elemSize)}; auto alloc = builder.create( loc, @@ -1603,7 +1603,7 @@ MLIRScanner::EmitGPUCallExpr(clang::CallExpr *expr) { args); ValueCategory(dst, /*isReference*/ true) .store(builder, - builder.create(loc, alloc, mt)); + builder.create(loc, mt, alloc)); auto retTy = getMLIRType(expr->getType()); return make_pair( ValueCategory(builder.create(loc, 0, retTy), @@ -1618,37 +1618,33 @@ MLIRScanner::EmitGPUCallExpr(clang::CallExpr *expr) { auto createBlockIdOp = [&](gpu::Dimension str, mlir::Type mlirType) -> mlir::Value { return builder.create( - loc, + loc, mlirType, builder.create( - loc, mlir::IndexType::get(builder.getContext()), str), - mlirType); + loc, mlir::IndexType::get(builder.getContext()), str)); }; auto createBlockDimOp = [&](gpu::Dimension str, mlir::Type mlirType) -> mlir::Value { return builder.create( - loc, + loc, mlirType, builder.create( - loc, mlir::IndexType::get(builder.getContext()), str), - mlirType); + loc, mlir::IndexType::get(builder.getContext()), str)); }; auto createThreadIdOp = [&](gpu::Dimension str, mlir::Type mlirType) -> mlir::Value { return builder.create( - loc, + loc, mlirType, builder.create( - loc, mlir::IndexType::get(builder.getContext()), str), - mlirType); + loc, mlir::IndexType::get(builder.getContext()), str)); }; auto createGridDimOp = [&](gpu::Dimension str, mlir::Type mlirType) -> mlir::Value { return builder.create( - loc, + loc, mlirType, builder.create( - loc, mlir::IndexType::get(builder.getContext()), str), - mlirType); + loc, mlir::IndexType::get(builder.getContext()), str)); }; if (auto ME = dyn_cast(ic->getSubExpr())) { @@ -1806,7 +1802,7 @@ ValueCategory MLIRScanner::VisitUnaryOperator(clang::UnaryOperator *U) { auto postTy = getMLIRType(U->getType()).cast(); if (postTy.getWidth() > 1) - res = builder.create(loc, res, postTy); + res = builder.create(loc, postTy, res); return ValueCategory(res, /*isReference*/ false); } case clang::UnaryOperator::Opcode::UO_Not: { @@ -1842,7 +1838,7 @@ ValueCategory MLIRScanner::VisitUnaryOperator(clang::UnaryOperator *U) { auto mt0 = mlir::MemRefType::get(shape, mt.getElementType(), MemRefLayoutAttrInterface(), mt.getMemorySpace()); - res = builder.create(loc, sub.val, mt0); + res = builder.create(loc, mt0, sub.val); return ValueCategory(res, /*isReference*/ false); } @@ -2026,7 +2022,7 @@ MLIRScanner::VisitUnaryExprOrTypeTraitExpr(UnaryExprOrTypeTraitExpr *Uop) { case UETT_SizeOf: { auto value = getTypeSize(Uop->getTypeOfArgument()); auto retTy = getMLIRType(Uop->getType()).cast(); - return ValueCategory(builder.create(loc, value, retTy), + return ValueCategory(builder.create(loc, retTy, value), /*isReference*/ false); } default: @@ -2131,9 +2127,9 @@ ValueCategory MLIRScanner::VisitBinaryOperator(clang::BinaryOperator *BO) { } if (postTy != prevTy) { if (signedType) { - res = builder.create(loc, res, postTy); + res = builder.create(loc, postTy, res); } else { - res = builder.create(loc, res, postTy); + res = builder.create(loc, postTy, res); } } return ValueCategory(res, /*isReference*/ false); @@ -2253,9 +2249,9 @@ ValueCategory MLIRScanner::VisitBinaryOperator(clang::BinaryOperator *BO) { auto prevTy = rhsv.getType().cast(); auto postTy = lhsv.getType().cast(); if (prevTy.getWidth() < postTy.getWidth()) - rhsv = builder.create(loc, rhsv, postTy); + rhsv = builder.create(loc, postTy, rhsv); if (prevTy.getWidth() > postTy.getWidth()) - rhsv = builder.create(loc, rhsv, postTy); + rhsv = builder.create(loc, postTy, rhsv); assert(lhsv.getType() == rhsv.getType()); if (signedType) return ValueCategory(builder.create(loc, lhsv, rhsv), @@ -2270,9 +2266,9 @@ ValueCategory MLIRScanner::VisitBinaryOperator(clang::BinaryOperator *BO) { auto prevTy = rhsv.getType().cast(); auto postTy = lhsv.getType().cast(); if (prevTy.getWidth() < postTy.getWidth()) - rhsv = builder.create(loc, rhsv, postTy); + rhsv = builder.create(loc, postTy, rhsv); if (prevTy.getWidth() > postTy.getWidth()) - rhsv = builder.create(loc, rhsv, postTy); + rhsv = builder.create(loc, postTy, rhsv); assert(lhsv.getType() == rhsv.getType()); return ValueCategory(builder.create(loc, lhsv, rhsv), /*isReference*/ false); @@ -2518,12 +2514,12 @@ ValueCategory MLIRScanner::VisitBinaryOperator(clang::BinaryOperator *BO) { if (prevTy.getWidth() < postTy.getWidth()) { if (signedType) { - tostore = builder.create(loc, tostore, postTy); + tostore = builder.create(loc, postTy, tostore); } else { - tostore = builder.create(loc, tostore, postTy); + tostore = builder.create(loc, postTy, tostore); } } else if (prevTy.getWidth() > postTy.getWidth()) { - tostore = builder.create(loc, tostore, postTy); + tostore = builder.create(loc, postTy, tostore); } } } @@ -2546,9 +2542,9 @@ ValueCategory MLIRScanner::VisitBinaryOperator(clang::BinaryOperator *BO) { auto prevTy = rhsV.getType().cast(); if (prevTy == postTy) { } else if (prevTy.getWidth() < postTy.getWidth()) { - rhsV = builder.create(loc, rhsV, postTy); + rhsV = builder.create(loc, postTy, rhsV); } else { - rhsV = builder.create(loc, rhsV, postTy); + rhsV = builder.create(loc, postTy, rhsV); } assert(rhsV.getType() == prev.getType()); result = builder.create(loc, prev, rhsV); @@ -2563,12 +2559,12 @@ ValueCategory MLIRScanner::VisitBinaryOperator(clang::BinaryOperator *BO) { if (prevTy == postTy) { } else if (prevTy.getWidth() < postTy.getWidth()) { if (signedType) { - rhsV = builder.create(loc, rhsV, postTy); + rhsV = builder.create(loc, postTy, rhsV); } else { - rhsV = builder.create(loc, rhsV, postTy); + rhsV = builder.create(loc, postTy, rhsV); } } else { - rhsV = builder.create(loc, rhsV, postTy); + rhsV = builder.create(loc, postTy, rhsV); } assert(rhsV.getType() == prev.getType()); result = builder.create(loc, prev, rhsV); @@ -2588,9 +2584,9 @@ ValueCategory MLIRScanner::VisitBinaryOperator(clang::BinaryOperator *BO) { auto postTy = getMLIRType(BO->getType()).cast(); if (prevTy.getWidth() < postTy.getWidth()) { - right = builder.create(loc, right, postTy); + right = builder.create(loc, postTy, right); } else { - right = builder.create(loc, right, postTy); + right = builder.create(loc, postTy, right); } } if (right.getType() != prev.getType()) { @@ -2617,9 +2613,9 @@ ValueCategory MLIRScanner::VisitBinaryOperator(clang::BinaryOperator *BO) { auto postTy = getMLIRType(BO->getType()).cast(); if (prevTy.getWidth() < postTy.getWidth()) { - right = builder.create(loc, right, postTy); + right = builder.create(loc, postTy, right); } else { - right = builder.create(loc, right, postTy); + right = builder.create(loc, postTy, right); } } if (right.getType() != prev.getType()) { @@ -2645,16 +2641,18 @@ ValueCategory MLIRScanner::VisitBinaryOperator(clang::BinaryOperator *BO) { auto postTy = prev.getType().cast(); if (prevTy.getWidth() < postTy.getWidth()) { - val = builder.create(loc, val, postTy); + val = builder.create(loc, postTy, val); } else if (prevTy.getWidth() > postTy.getWidth()) { - val = builder.create(loc, val, postTy); + val = builder.create(loc, postTy, val); } - result = builder.create(loc, prev, val); + result = builder.create(loc, prev, val); } else { if (signedType) - result = builder.create(loc, prev, rhs.getValue(builder)); + result = + builder.create(loc, prev, rhs.getValue(builder)); else - result = builder.create(loc, prev, rhs.getValue(builder)); + result = + builder.create(loc, prev, rhs.getValue(builder)); } lhs.store(builder, result); return lhs; @@ -3055,7 +3053,7 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { auto ty = mlir::MemRefType::get(mt.getShape(), mt.getElementType(), MemRefLayoutAttrInterface(), ut.getMemorySpace()); - return ValueCategory(builder.create(loc, se.val, ty), + return ValueCategory(builder.create(loc, ty, se.val), /*isReference*/ se.isReference); } case clang::CastKind::CK_BitCast: { @@ -3094,14 +3092,14 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { E->getType()->getUnqualifiedDesugaredType()) ->getPointeeType()); mlir::Value allocSize = builder.create( - loc, Visit(CI->getArg(0)).getValue(builder), - mlir::IndexType::get(builder.getContext())); + loc, mlir::IndexType::get(builder.getContext()), + Visit(CI->getArg(0)).getValue(builder)); if (sr->getDecl()->getName() == "calloc") { allocSize = builder.create( loc, allocSize, - builder.create( - loc, Visit(CI->getArg(1)).getValue(builder), - mlir::IndexType::get(builder.getContext()))); + builder.create( + loc, mlir::IndexType::get(builder.getContext()), + Visit(CI->getArg(1)).getValue(builder))); } mlir::Value args[1] = { builder.create(loc, allocSize, elemSize)}; @@ -3122,10 +3120,11 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { .getAddressSpace()), val); } - auto i8_0 = builder.create(loc, 0, 8); - auto sizev = builder.create(loc, allocSize, - builder.getI64Type()); - auto falsev = builder.create(loc, false, 1); + auto i8_0 = builder.create(loc, 0, 8); + auto sizev = builder.create( + loc, builder.getI64Type(), allocSize); + auto falsev = + builder.create(loc, false, 1); builder.create(loc, val, i8_0, sizev, falsev); } return ValueCategory(alloc, /*isReference*/ false); @@ -3179,9 +3178,8 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { scalar)), /*isReference*/ false); } - return ValueCategory( - builder.create(loc, scalar, ty), - /*isReference*/ false); + return ValueCategory(builder.create(loc, ty, scalar), + /*isReference*/ false); } else { E->dump(); E->getType()->dump(); @@ -3233,11 +3231,13 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { signedType = true; } if (signedType) - return ValueCategory(builder.create(loc, scalar, ty), - /*isReference*/ false); + return ValueCategory( + builder.create(loc, ty, scalar), + /*isReference*/ false); else - return ValueCategory(builder.create(loc, scalar, ty), - /*isReference*/ false); + return ValueCategory( + builder.create(loc, ty, scalar), + /*isReference*/ false); } case clang::CastKind::CK_FloatingToIntegral: { auto scalar = Visit(E->getSubExpr()).getValue(builder); @@ -3250,11 +3250,13 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { signedType = true; } if (signedType) - return ValueCategory(builder.create(loc, scalar, ty), - /*isReference*/ false); + return ValueCategory( + builder.create(loc, ty, scalar), + /*isReference*/ false); else - return ValueCategory(builder.create(loc, scalar, ty), - /*isReference*/ false); + return ValueCategory( + builder.create(loc, ty, scalar), + /*isReference*/ false); } case clang::CastKind::CK_IntegralCast: { auto scalar = Visit(E->getSubExpr()).getValue(builder); @@ -3267,7 +3269,7 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { } if (scalar.getType().isa() || postTy.isa()) { - return ValueCategory(builder.create(loc, scalar, postTy), + return ValueCategory(builder.create(loc, postTy, scalar), false); } if (!scalar.getType().isa()) { @@ -3287,7 +3289,7 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { return ValueCategory(scalar, /*isReference*/ false); if (prevTy.getWidth() < postTy.getWidth()) { if (signedType) { - if (auto CI = scalar.getDefiningOp()) { + if (auto CI = scalar.getDefiningOp()) { return ValueCategory( builder.create( loc, postTy, @@ -3297,10 +3299,10 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { /*isReference*/ false); } return ValueCategory( - builder.create(loc, scalar, postTy), + builder.create(loc, postTy, scalar), /*isReference*/ false); } else { - if (auto CI = scalar.getDefiningOp()) { + if (auto CI = scalar.getDefiningOp()) { return ValueCategory( builder.create( loc, postTy, @@ -3310,7 +3312,7 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { /*isReference*/ false); } return ValueCategory( - builder.create(loc, scalar, postTy), + builder.create(loc, postTy, scalar), /*isReference*/ false); } } else { @@ -3323,7 +3325,7 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { postTy.getWidth()))), /*isReference*/ false); } - return ValueCategory(builder.create(loc, scalar, postTy), + return ValueCategory(builder.create(loc, postTy, scalar), /*isReference*/ false); } } @@ -3348,10 +3350,10 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { false); } if (prevTy.getWidth() < postTy.getWidth()) { - return ValueCategory(builder.create(loc, scalar, postTy), + return ValueCategory(builder.create(loc, postTy, scalar), /*isReference*/ false); } else { - return ValueCategory(builder.create(loc, scalar, postTy), + return ValueCategory(builder.create(loc, postTy, scalar), /*isReference*/ false); } } @@ -3440,9 +3442,9 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { } if (postTy.getWidth() > 1) { if (signedType) { - res = builder.create(loc, res, postTy); + res = builder.create(loc, postTy, res); } else { - res = builder.create(loc, res, postTy); + res = builder.create(loc, postTy, res); } } return ValueCategory(res, /*isReference*/ false); @@ -3463,9 +3465,9 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) { res = builder.create(loc, CmpFPredicate::UNE, res, Zero); if (1 < postTy.getWidth()) { if (signedType) { - res = builder.create(loc, res, postTy); + res = builder.create(loc, postTy, res); } else { - res = builder.create(loc, res, postTy); + res = builder.create(loc, postTy, res); } } return ValueCategory(res, /*isReference*/ false); diff --git a/tools/mlir-clang/Lib/pragmaHandler.cc b/tools/mlir-clang/Lib/pragmaHandler.cc index 30ff1e748c5c..c4bd663d4746 100644 --- a/tools/mlir-clang/Lib/pragmaHandler.cc +++ b/tools/mlir-clang/Lib/pragmaHandler.cc @@ -127,7 +127,7 @@ class PragmaLowerToHandler : public PragmaHandler { /// the function decl that lower_to is attached to with that MLIR function /// symbol in the class-referenced dictionary. /// - // #pragma lower_to(copy_op, "linalg.copy") input(a), output(b) + // #pragma lower_to(copy_op, "memref.copy") input(a), output(b) void HandlePragma(Preprocessor &PP, PragmaIntroducer Introducer, Token &PragmaTok) override { Token Tok; diff --git a/tools/mlir-clang/Lib/utils.cc b/tools/mlir-clang/Lib/utils.cc index f46f3664cba6..1595333d6633 100644 --- a/tools/mlir-clang/Lib/utils.cc +++ b/tools/mlir-clang/Lib/utils.cc @@ -31,10 +31,10 @@ using namespace clang; Operation *buildLinalgOp(StringRef name, OpBuilder &b, SmallVectorImpl &input, SmallVectorImpl &output) { - if (name.compare("linalg.copy") == 0) { - assert(input.size() == 1 && "linalg::copyOp requires 1 input"); - assert(output.size() == 1 && "linalg::CopyOp requires 1 output"); - return b.create(b.getUnknownLoc(), input[0], output[0]); + if (name.compare("memref.copy") == 0) { + assert(input.size() == 1 && "memref::copyOp requires 1 input"); + assert(output.size() == 1 && "memref::CopyOp requires 1 output"); + return b.create(b.getUnknownLoc(), input[0], output[0]); } else { llvm::report_fatal_error(llvm::Twine("builder not supported for: ") + name); return nullptr; @@ -49,7 +49,7 @@ mlirclang::replaceFuncByOperation(FuncOp f, StringRef opName, OpBuilder &b, assert(ctx->isOperationRegistered(opName) && "Provided lower_to opName should be registered."); - if (opName.startswith("linalg")) + if (opName.startswith("memref")) return buildLinalgOp(opName, b, input, output); // NOTE: The attributes of the provided FuncOp is ignored. diff --git a/tools/mlir-clang/Test/Verification/charswitch.cpp b/tools/mlir-clang/Test/Verification/charswitch.cpp index f70f56f8e183..6ff1c6e784ec 100644 --- a/tools/mlir-clang/Test/Verification/charswitch.cpp +++ b/tools/mlir-clang/Test/Verification/charswitch.cpp @@ -33,6 +33,6 @@ int foo(char t) { // CHECK-NEXT: 65: ^bb1(%c30_i32, %true, %0 : i32, i1, i32) // CHECK-NEXT: ] // CHECK-NEXT: ^bb1(%2: i32, %3: i1, %4: i32): // 3 preds: ^bb0, ^bb0, ^bb0 -// CHECK-NEXT: %5 = select %3, %2, %4 : i32 +// CHECK-NEXT: %5 = arith.select %3, %2, %4 : i32 // CHECK-NEXT: return %5 : i32 // CHECK-NEXT: } diff --git a/tools/mlir-clang/Test/Verification/combif.c b/tools/mlir-clang/Test/Verification/combif.c index 89016bc45721..8d77f9994e9a 100644 --- a/tools/mlir-clang/Test/Verification/combif.c +++ b/tools/mlir-clang/Test/Verification/combif.c @@ -34,7 +34,7 @@ int solver( float** y, // CHECK-NEXT: %3 = arith.andi %2, %arg6 : i1 // CHECK-NEXT: scf.condition(%3) %arg4, %arg5 : f32, i32 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg4: f32, %arg5: i32): // no predecessors +// CHECK-NEXT: ^bb0(%arg4: f32, %arg5: i32): // CHECK-NEXT: %2 = arith.cmpf ugt, %arg3, %cst : f32 // CHECK-NEXT: %3:2 = scf.if %2 -> (f32, i1) { // CHECK-NEXT: scf.yield %arg4, %false : f32, i1 diff --git a/tools/mlir-clang/Test/Verification/cond.c b/tools/mlir-clang/Test/Verification/cond.c index 7e31d2f1f3c1..8cad47002fe8 100644 --- a/tools/mlir-clang/Test/Verification/cond.c +++ b/tools/mlir-clang/Test/Verification/cond.c @@ -23,6 +23,6 @@ int set (int b) // CHCEK-NEXT: %c1_i32 = constant 1 : i32 // CHCEK-NEXT: %c2_i32 = constant 2 : i32 // CHCEK-NEXT: %0 = arith.trunci %arg0 : i32 to i1 -// CHCEK-NEXT: %1 = select %0, %c1_i32, %c2_i32 : i32 +// CHCEK-NEXT: %1 = arith.select %0, %c1_i32, %c2_i32 : i32 // CHCEK-NEXT: return %1 : i32 // CHCEK-NEXT: } diff --git a/tools/mlir-clang/Test/Verification/cond2.c b/tools/mlir-clang/Test/Verification/cond2.c index dd271fb0272e..e64d30b58e5e 100644 --- a/tools/mlir-clang/Test/Verification/cond2.c +++ b/tools/mlir-clang/Test/Verification/cond2.c @@ -21,6 +21,6 @@ int set (int b) // CHCEK-NEXT: %c1_i32 = constant 1 : i32 // CHCEK-NEXT: %c2_i32 = constant 2 : i32 // CHCEK-NEXT: %0 = arith.trunci %arg0 : i32 to i1 -// CHCEK-NEXT: %1 = select %0, %c2_i32, %c1_i32 : i32 +// CHCEK-NEXT: %1 = arith.select %0, %c2_i32, %c1_i32 : i32 // CHCEK-NEXT: return %1 : i32 // CHCEK-NEXT: } diff --git a/tools/mlir-clang/Test/Verification/constexpr.cpp b/tools/mlir-clang/Test/Verification/constexpr.cpp index 8ed6f5dce04f..f523e5173709 100644 --- a/tools/mlir-clang/Test/Verification/constexpr.cpp +++ b/tools/mlir-clang/Test/Verification/constexpr.cpp @@ -21,7 +21,7 @@ int foo() { // CHECK-NEXT: %6 = arith.cmpi slt, %arg0, %5 : i32 // CHECK-NEXT: scf.condition(%6) %arg0 : i32 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg0: i32): // no predecessors +// CHECK-NEXT: ^bb0(%arg0: i32): // CHECK-NEXT: %5 = arith.index_cast %arg0 : i32 to index // CHECK-NEXT: memref.store %arg0, %0[%5] : memref<14xi32> // CHECK-NEXT: %6 = arith.addi %arg0, %c1_i32 : i32 diff --git a/tools/mlir-clang/Test/Verification/gcd.c b/tools/mlir-clang/Test/Verification/gcd.c index 2696677c7023..53692ed38999 100644 --- a/tools/mlir-clang/Test/Verification/gcd.c +++ b/tools/mlir-clang/Test/Verification/gcd.c @@ -15,7 +15,7 @@ int gcd(int m, int n) { // CHECK-NEXT: %1 = arith.cmpi sgt, %arg2, %c0_i32 : i32 // CHECK-NEXT: scf.condition(%1) %arg3, %arg2 : i32, i32 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg2: i32, %arg3: i32): // no predecessors +// CHECK-NEXT: ^bb0(%arg2: i32, %arg3: i32): // CHECK-NEXT: %1 = arith.remsi %arg2, %arg3 : i32 // CHECK-NEXT: scf.yield %1, %arg3 : i32, i32 // CHECK-NEXT: } diff --git a/tools/mlir-clang/Test/Verification/loop.cpp b/tools/mlir-clang/Test/Verification/loop.cpp index ecc2a0dd6e4a..b4012d9c95be 100644 --- a/tools/mlir-clang/Test/Verification/loop.cpp +++ b/tools/mlir-clang/Test/Verification/loop.cpp @@ -26,7 +26,7 @@ void div_(int* sizes) { // CHECK-NEXT: %5 = arith.cmpi slt, %arg1, %4 : i32 // CHECK-NEXT: scf.condition(%5) %arg1 : i32 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg1: i32): // no predecessors +// CHECK-NEXT: ^bb0(%arg1: i32): // CHECK-NEXT: %4 = arith.index_cast %arg1 : i32 to index // CHECK-NEXT: %5 = arith.index_cast %4 : index to i64 // CHECK-NEXT: %6 = llvm.getelementptr %2[%5] : (!llvm.ptr>, i64) -> !llvm.ptr> diff --git a/tools/mlir-clang/Test/Verification/lower-to-linalg-op.c b/tools/mlir-clang/Test/Verification/lower-to-linalg-op.c index 87d349c91766..ef0fd86d5223 100644 --- a/tools/mlir-clang/Test/Verification/lower-to-linalg-op.c +++ b/tools/mlir-clang/Test/Verification/lower-to-linalg-op.c @@ -1,6 +1,6 @@ // RUN: mlir-clang %s -S | FileCheck %s -#pragma lower_to(copy_op, "linalg.copy") "input"(a), "output"(b) +#pragma lower_to(copy_op, "memref.copy") "input"(a), "output"(b) void copy_op(int b[3][3], int a[3][3]) { for (int i = 0; i < 3; i++) for (int j = 0; j < 3; j++) @@ -10,7 +10,7 @@ void copy_op(int b[3][3], int a[3][3]) { int main() { int a[3][3]; int b[3][3]; - // CHECK: linalg.copy({{.*}}, {{.*}}) : memref<3x3xi32>, memref<3x3xi32> + // CHECK: memref.copy {{.*}}, {{.*}} : memref<3x3xi32> to memref<3x3xi32> copy_op(a, b); return 0; } diff --git a/tools/mlir-clang/Test/Verification/min.c b/tools/mlir-clang/Test/Verification/min.c index 3c9dc260c70b..d7580915bdb9 100644 --- a/tools/mlir-clang/Test/Verification/min.c +++ b/tools/mlir-clang/Test/Verification/min.c @@ -10,7 +10,7 @@ int min(int a, int b) { // CHECK-NEXT: %0 = llvm.mlir.undef : i32 // CHECK-NEXT: %1 = arith.cmpi slt, %arg0, %arg1 : i32 // CHECK-NEXT: %2 = arith.cmpi sge, %arg0, %arg1 : i32 -// CHECK-NEXT: %3 = select %1, %arg0, %0 : i32 -// CHECK-NEXT: %4 = select %2, %arg1, %3 : i32 +// CHECK-NEXT: %3 = arith.select %1, %arg0, %0 : i32 +// CHECK-NEXT: %4 = arith.select %2, %arg1, %3 : i32 // CHECK-NEXT: return %4 : i32 // CHECK-NEXT: } diff --git a/tools/mlir-clang/Test/Verification/nus.c b/tools/mlir-clang/Test/Verification/nus.c index 7b2d0310eefe..9e3888b77ab6 100644 --- a/tools/mlir-clang/Test/Verification/nus.c +++ b/tools/mlir-clang/Test/Verification/nus.c @@ -10,7 +10,7 @@ // CHECK-NEXT: %0 = affine.load %arg1[%arg2] : memref // CHECK-NEXT: %1 = affine.load %arg1[%arg2 - 1] : memref // CHECK-NEXT: %2 = arith.cmpi sge, %0, %1 : i32 -// CHECK-NEXT: %3 = select %2, %0, %1 : i32 +// CHECK-NEXT: %3 = arith.select %2, %0, %1 : i32 // CHECK-NEXT: affine.store %3, %arg1[%arg2] : memref // CHECK-NEXT: } // CHECK-NEXT: } diff --git a/tools/mlir-clang/Test/Verification/scor.c b/tools/mlir-clang/Test/Verification/scor.c index 6cd5093f9673..b8ab1c8a825d 100644 --- a/tools/mlir-clang/Test/Verification/scor.c +++ b/tools/mlir-clang/Test/Verification/scor.c @@ -32,7 +32,7 @@ void kernel_correlation(double out[28], double stddev[28], _Bool cmp) // CHECK: func @kernel_correlation(%arg0: memref, %arg1: memref, %arg2: i8) // CHECK-DAG: %[[cst:.+]] = arith.constant 1.000000e+00 : f64 // CHECK-DAG: %[[cst_0:.+]] = arith.constant 3.140000e+00 : f64 -// CHECK: %[[sel:.+]] = select %{{.*}}, %[[cst]], %[[cst_0]] : f64 +// CHECK: %[[sel:.+]] = arith.select %{{.*}}, %[[cst]], %[[cst_0]] : f64 // CHECK-NEXT: affine.for %arg3 = 0 to 28 { // CHECK-NEXT: affine.store %[[cst_0]], %arg1[%arg3] : memref // CHECK-NEXT: affine.store %[[sel]], %arg0[%arg3] : memref diff --git a/tools/mlir-clang/Test/Verification/sec.c b/tools/mlir-clang/Test/Verification/sec.c index de3bd9e233aa..57f93981d6d9 100644 --- a/tools/mlir-clang/Test/Verification/sec.c +++ b/tools/mlir-clang/Test/Verification/sec.c @@ -22,7 +22,7 @@ int compute_tran_temp(int total_iterations, int num_iterations) // CHECK-NEXT: %3 = arith.cmpf ult, %arg4, %0 : f32 // CHECK-NEXT: scf.condition(%3) %arg2, %arg3, %arg4 : i32, i32, f32 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg2: i32, %arg3: i32, %arg4: f32): // no predecessors +// CHECK-NEXT: ^bb0(%arg2: i32, %arg3: i32, %arg4: f32): // CHECK-NEXT: %3 = arith.addf %arg4, %1 : f32 // CHECK-NEXT: scf.yield %arg3, %arg2, %3 : i32, i32, f32 // CHECK-NEXT: } diff --git a/tools/mlir-clang/Test/Verification/snus.c b/tools/mlir-clang/Test/Verification/snus.c index f81f0193622b..c60bb1330e21 100644 --- a/tools/mlir-clang/Test/Verification/snus.c +++ b/tools/mlir-clang/Test/Verification/snus.c @@ -12,7 +12,6 @@ void kernel_nussinov(double* out, int n) { for (int k=0; k<10; k++) { out[n] = max_score(out[n], table[k]); } - //} #pragma endscop } @@ -25,7 +24,7 @@ void kernel_nussinov(double* out, int n) { // CHECK-NEXT: %4 = affine.for %arg2 = 0 to 10 iter_args(%arg3 = %3) -> (f64) { // CHECK-NEXT: %5 = affine.load %[[i0]][%arg2] : memref<20xf64> // CHECK-NEXT: %6 = arith.cmpf uge, %arg3, %5 : f64 -// CHECK-NEXT: %7 = select %6, %arg3, %5 : f64 +// CHECK-NEXT: %7 = arith.select %6, %arg3, %5 : f64 // CHECK-NEXT: affine.yield %7 : f64 // CHECK-NEXT: } // CHECK-NEXT: affine.store %4, %arg0[symbol(%[[i1]])] : memref diff --git a/tools/mlir-clang/Test/Verification/switcherr.c b/tools/mlir-clang/Test/Verification/switcherr.c index 3e7de7c0e500..29f57b213ec9 100644 --- a/tools/mlir-clang/Test/Verification/switcherr.c +++ b/tools/mlir-clang/Test/Verification/switcherr.c @@ -30,6 +30,6 @@ int foo(int t) { // CHECK-NEXT: 2: ^bb1(%c30_i32, %true, %0 : i32, i1, i32) // CHECK-NEXT: ] // CHECK-NEXT: ^bb1(%1: i32, %2: i1, %3: i32): // 3 preds: ^bb0, ^bb0, ^bb0 -// CHECK-NEXT: %4 = select %2, %1, %3 : i32 +// CHECK-NEXT: %4 = arith.select %2, %1, %3 : i32 // CHECK-NEXT: return %4 : i32 // CHECK-NEXT: } diff --git a/tools/mlir-clang/Test/Verification/switchnone.c b/tools/mlir-clang/Test/Verification/switchnone.c index 0446b605e97d..bd6a3a45ad97 100644 --- a/tools/mlir-clang/Test/Verification/switchnone.c +++ b/tools/mlir-clang/Test/Verification/switchnone.c @@ -6,7 +6,6 @@ int foo(int t) { return t; } -// TODO the select should be canonicalized better // CHECK: func @foo(%arg0: i32) -> i32 attributes {llvm.linkage = #llvm.linkage} { // CHECK-NEXT: return %arg0 : i32 // CHECK-NEXT: } diff --git a/tools/mlir-clang/Test/Verification/whileset.c b/tools/mlir-clang/Test/Verification/whileset.c index 3edcc98b2ffb..5073099420be 100644 --- a/tools/mlir-clang/Test/Verification/whileset.c +++ b/tools/mlir-clang/Test/Verification/whileset.c @@ -30,7 +30,7 @@ void set (int path[20]) // CHECK-NEXT: %0 = scf.while (%arg1 = %c0_i32, %arg2 = %true) : (i32, i1) -> i32 { // CHECK-NEXT: scf.condition(%arg2) %arg1 : i32 // CHECK-NEXT: } do { -// CHECK-NEXT: ^bb0(%arg1: i32): // no predecessors +// CHECK-NEXT: ^bb0(%arg1: i32): // CHECK-NEXT: %1 = arith.index_cast %arg1 : i32 to index // CHECK-NEXT: memref.store %c3_i32, %arg0[%1] : memref // CHECK-NEXT: %2 = arith.addi %arg1, %c1_i32 : i32 diff --git a/tools/mlir-clang/Test/polybench/medley/floyd-warshall/floyd-warshall.c b/tools/mlir-clang/Test/polybench/medley/floyd-warshall/floyd-warshall.c index ec193a20c986..dc7361db7794 100644 --- a/tools/mlir-clang/Test/polybench/medley/floyd-warshall/floyd-warshall.c +++ b/tools/mlir-clang/Test/polybench/medley/floyd-warshall/floyd-warshall.c @@ -136,7 +136,7 @@ int main(int argc, char** argv) // CHECK-NEXT: %3 = affine.load %arg1[%arg2, %arg4] : memref // CHECK-NEXT: %4 = arith.addi %2, %3 : i32 // CHECK-NEXT: %5 = arith.cmpi slt, %1, %4 : i32 -// CHECK-NEXT: %6 = select %5, %1, %4 : i32 +// CHECK-NEXT: %6 = arith.select %5, %1, %4 : i32 // CHECK-NEXT: affine.store %6, %arg1[%arg3, %arg4] : memref // CHECK-NEXT: } // CHECK-NEXT: } diff --git a/tools/mlir-clang/Test/polybench/stencils/jacobi-1d/jacobi-1d.c b/tools/mlir-clang/Test/polybench/stencils/jacobi-1d/jacobi-1d.c index 570c47e60512..098e3c474fd0 100644 --- a/tools/mlir-clang/Test/polybench/stencils/jacobi-1d/jacobi-1d.c +++ b/tools/mlir-clang/Test/polybench/stencils/jacobi-1d/jacobi-1d.c @@ -144,7 +144,7 @@ int main(int argc, char** argv) // CHECK-NEXT: %4 = arith.addf %2, %3 : f64 // CHECK-NEXT: %5 = affine.load %arg2[%arg5 + 1] : memref // CHECK-NEXT: %6 = arith.addf %4, %5 : f64 -// CHECK-NEXT: %7 = arith.mulf %cst, %6 : f64 +// CHECK-NEXT: %7 = arith.mulf %6, %cst : f64 // CHECK-NEXT: affine.store %7, %arg3[%arg5] : memref // CHECK-NEXT: } // CHECK-NEXT: affine.for %arg5 = 1 to #map()[%0] { @@ -153,7 +153,7 @@ int main(int argc, char** argv) // CHECK-NEXT: %4 = arith.addf %2, %3 : f64 // CHECK-NEXT: %5 = affine.load %arg3[%arg5 + 1] : memref // CHECK-NEXT: %6 = arith.addf %4, %5 : f64 -// CHECK-NEXT: %7 = arith.mulf %cst, %6 : f64 +// CHECK-NEXT: %7 = arith.mulf %6, %cst : f64 // CHECK-NEXT: affine.store %7, %arg2[%arg5] : memref // CHECK-NEXT: } // CHECK-NEXT: } diff --git a/tools/mlir-clang/Test/polybench/stencils/jacobi-2d/jacobi-2d.c b/tools/mlir-clang/Test/polybench/stencils/jacobi-2d/jacobi-2d.c index 58b0d8e696d4..9ddb349a9564 100644 --- a/tools/mlir-clang/Test/polybench/stencils/jacobi-2d/jacobi-2d.c +++ b/tools/mlir-clang/Test/polybench/stencils/jacobi-2d/jacobi-2d.c @@ -150,7 +150,7 @@ int main(int argc, char** argv) // CHECK-NEXT: %8 = arith.addf %6, %7 : f64 // CHECK-NEXT: %9 = affine.load %arg2[%arg5 - 1, %arg6] : memref // CHECK-NEXT: %10 = arith.addf %8, %9 : f64 -// CHECK-NEXT: %11 = arith.mulf %cst, %10 : f64 +// CHECK-NEXT: %11 = arith.mulf %10, %cst : f64 // CHECK-NEXT: affine.store %11, %arg3[%arg5, %arg6] : memref // CHECK-NEXT: } // CHECK-NEXT: } @@ -165,7 +165,7 @@ int main(int argc, char** argv) // CHECK-NEXT: %8 = arith.addf %6, %7 : f64 // CHECK-NEXT: %9 = affine.load %arg3[%arg5 - 1, %arg6] : memref // CHECK-NEXT: %10 = arith.addf %8, %9 : f64 -// CHECK-NEXT: %11 = arith.mulf %cst, %10 : f64 +// CHECK-NEXT: %11 = arith.mulf %10, %cst : f64 // CHECK-NEXT: affine.store %11, %arg2[%arg5, %arg6] : memref // CHECK-NEXT: } // CHECK-NEXT: } diff --git a/tools/mlir-clang/mlir-clang.cc b/tools/mlir-clang/mlir-clang.cc index da67a9adb5b7..8583029200be 100644 --- a/tools/mlir-clang/mlir-clang.cc +++ b/tools/mlir-clang/mlir-clang.cc @@ -23,8 +23,8 @@ #include "mlir/Conversion/LLVMCommon/LoweringOptions.h" #include "mlir/Conversion/MathToLLVM/MathToLLVM.h" #include "mlir/Conversion/OpenMPToLLVM/ConvertOpenMPToLLVM.h" +#include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h" #include "mlir/Conversion/SCFToOpenMP/SCFToOpenMP.h" -#include "mlir/Conversion/SCFToStandard/SCFToStandard.h" #include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVMPass.h" #include "mlir/Dialect/Affine/Passes.h" #include "mlir/Dialect/GPU/GPUDialect.h" @@ -97,10 +97,6 @@ static cl::opt CUDAPath("cuda-path", cl::init(""), static cl::opt Output("o", cl::init("-"), cl::desc("Output file")); -static cl::list inputFileName(cl::Positional, cl::OneOrMore, - cl::desc(""), - cl::cat(toolOptions)); - static cl::opt cfunction("function", cl::desc(""), cl::init("main"), cl::cat(toolOptions)); @@ -370,6 +366,10 @@ int main(int argc, char **argv) { } using namespace mlir; + cl::list inputFileName(cl::Positional, cl::OneOrMore, + cl::desc(""), + cl::cat(toolOptions)); + int size = MLIRArgs.size(); const char **data = MLIRArgs.data(); InitLLVM y(size, data); @@ -383,8 +383,6 @@ int main(int argc, char **argv) { } } - // registerDialect(); - // registerDialect(); mlir::DialectRegistry registry; mlir::registerOpenMPDialectTranslation(registry); mlir::registerLLVMDialectTranslation(registry); @@ -402,7 +400,6 @@ int main(int argc, char **argv) { context.getOrLoadDialect(); context.getOrLoadDialect(); context.getOrLoadDialect(); - // MLIRContext context; LLVM::LLVMPointerType::attachInterface(context); LLVM::LLVMStructType::attachInterface(context); @@ -577,7 +574,7 @@ int main(int argc, char **argv) { } module->walk([&](mlir::omp::ParallelOp) { LinkOMP = true; }); mlir::PassManager pm3(&context); - pm3.addPass(mlir::createLowerToCFGPass()); + pm3.addPass(mlir::createConvertSCFToCFPass()); LowerToLLVMOptions options(&context); options.dataLayout = DL; // invalid for gemm.c init array