Bump LLVM to llvm/llvm-project@9261ab708e37 (#16195)
diff --git a/compiler/plugins/input/StableHLO/stablehlo-iree/Conversion/Preprocessing/Canonicalization.cpp b/compiler/plugins/input/StableHLO/stablehlo-iree/Conversion/Preprocessing/Canonicalization.cpp
index 160c048..66206e5 100644
--- a/compiler/plugins/input/StableHLO/stablehlo-iree/Conversion/Preprocessing/Canonicalization.cpp
+++ b/compiler/plugins/input/StableHLO/stablehlo-iree/Conversion/Preprocessing/Canonicalization.cpp
@@ -142,7 +142,7 @@
// The canonical form has the constant operand as the RHS.
if (isa<IntegerType>(type.getElementType()) && lhsAttr && !rhsAttr) {
- rewriter.updateRootInPlace(op, [op, lhs, rhs] {
+ rewriter.modifyOpInPlace(op, [op, lhs, rhs] {
op->setOperands(ValueRange{rhs, lhs});
});
return success();
@@ -239,7 +239,7 @@
// The canonical form has the constant operand as the RHS.
if (isa<IntegerType>(type.getElementType()) && lhsAttr && !rhsAttr) {
- rewriter.updateRootInPlace(op, [op, lhs, rhs] {
+ rewriter.modifyOpInPlace(op, [op, lhs, rhs] {
op->setOperands(ValueRange{rhs, lhs});
});
return success();
@@ -381,7 +381,7 @@
// The canonical form has the constant operand as the RHS.
if (lhsAttr && !rhsAttr) {
- rewriter.updateRootInPlace(op, [&op, direction, lhs, rhs] {
+ rewriter.modifyOpInPlace(op, [&op, direction, lhs, rhs] {
op.setComparisonDirection(invertDirection(direction));
op->setOperands(ValueRange{rhs, lhs});
});
@@ -1101,7 +1101,7 @@
int operandNum = operand.getOperandNumber();
auto emptyTensorOp = rewriter.create<tensor::EmptyOp>(
loc, operandType->getShape(), operandType->getElementType());
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
owner, [&]() { owner->setOperand(operandNum, emptyTensorOp); });
didUpdate = true;
}
diff --git a/compiler/plugins/input/StableHLO/stablehlo-iree/Conversion/StableHLOToIREEInputDialects.cpp b/compiler/plugins/input/StableHLO/stablehlo-iree/Conversion/StableHLOToIREEInputDialects.cpp
index f82fb9a..0f42f5b 100644
--- a/compiler/plugins/input/StableHLO/stablehlo-iree/Conversion/StableHLOToIREEInputDialects.cpp
+++ b/compiler/plugins/input/StableHLO/stablehlo-iree/Conversion/StableHLOToIREEInputDialects.cpp
@@ -318,7 +318,7 @@
convertedResultTypes);
// Update the function in place.
- rewriter.startRootUpdate(srcOp);
+ rewriter.startOpModification(srcOp);
srcOp.setType(newFuncType);
rewriteFuncAttrs(srcOp);
setFuncEncodings(srcOp, oldFuncType, newFuncType);
@@ -330,7 +330,7 @@
return failure();
}
- rewriter.finalizeRootUpdate(srcOp);
+ rewriter.finalizeOpModification(srcOp);
return success();
}
};
@@ -371,7 +371,7 @@
return rewriter.notifyMatchFailure(globalOp,
"result type conversion failed");
}
- rewriter.updateRootInPlace(globalOp, [&]() {
+ rewriter.modifyOpInPlace(globalOp, [&]() {
globalOp.setType(newType);
if (Attribute oldValue = globalOp.getValueAttr()) {
globalOp.setValueAttr(
diff --git a/compiler/plugins/input/StableHLO/stablehlo-iree/Conversion/test/stablehlo_to_linalg_reduce.mlir b/compiler/plugins/input/StableHLO/stablehlo-iree/Conversion/test/stablehlo_to_linalg_reduce.mlir
index 83d7ca6..de2fd82 100644
--- a/compiler/plugins/input/StableHLO/stablehlo-iree/Conversion/test/stablehlo_to_linalg_reduce.mlir
+++ b/compiler/plugins/input/StableHLO/stablehlo-iree/Conversion/test/stablehlo_to_linalg_reduce.mlir
@@ -33,7 +33,7 @@
// CHECK-PRIMITIVE-DAG: %[[INIT:.*]] = tensor.extract %{{.*}} : tensor<i32>
// CHECK-PRIMITIVE-DAG: %[[INIT_TENSOR:.*]] = tensor.empty()
// CHECK-PRIMITIVE-DAG: %[[FILL_TENSOR:.*]] = linalg.fill ins(%[[INIT]]{{.*}}outs(%[[INIT_TENSOR]]
-// CHECK-PRIMITIVE: linalg.reduce { arith.addi }
+// CHECK-PRIMITIVE: linalg.reduce { arith.addi {overflowFlags = #arith.overflow<none>} }
// CHECK-PRIMITIVE-SAME: ins(%{{.*}}tensor<5x4xi32>)
// CHECK-PRIMITIVE-SAME: outs(%[[FILL_TENSOR]] : tensor<5xi32>)
// CHECK-PRIMITIVE-SAME: dimensions = [1] {someattr}
diff --git a/compiler/src/iree/compiler/Codegen/Common/ConcretizePadResultShape.cpp b/compiler/src/iree/compiler/Codegen/Common/ConcretizePadResultShape.cpp
index bb0f42c..aac8334 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ConcretizePadResultShape.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ConcretizePadResultShape.cpp
@@ -120,8 +120,8 @@
staticShape, padOp.getResultType().getElementType(),
padOp.getResultType().getEncoding());
- rewriter.updateRootInPlace(
- padOp, [&]() { padOp.getResult().setType(resultType); });
+ rewriter.modifyOpInPlace(padOp,
+ [&]() { padOp.getResult().setType(resultType); });
return success();
}
};
diff --git a/compiler/src/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp b/compiler/src/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp
index 7eea1e0..d37b808 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp
@@ -511,7 +511,7 @@
LogicalResult matchAndRewrite(linalg::LinalgOp op,
PatternRewriter &rewriter) const override {
- rewriter.startRootUpdate(op);
+ rewriter.startOpModification(op);
bool modifiedOutput = false;
Location loc = op.getLoc();
for (OpOperand &opOperand : op.getDpsInitsMutable()) {
@@ -534,10 +534,10 @@
op->setOperand(opOperand.getOperandNumber(), fillOp);
}
if (!modifiedOutput) {
- rewriter.cancelRootUpdate(op);
+ rewriter.cancelOpModification(op);
return failure();
}
- rewriter.finalizeRootUpdate(op);
+ rewriter.finalizeOpModification(op);
return success();
}
};
diff --git a/compiler/src/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp b/compiler/src/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp
index 862b835..35399ce 100644
--- a/compiler/src/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp
@@ -327,8 +327,8 @@
return rewriter.notifyMatchFailure(op, "unhandled non-zero offset");
}
- rewriter.updateRootInPlace(op,
- [&] { op->setOperand(0, adaptor.getSource()); });
+ rewriter.modifyOpInPlace(op,
+ [&] { op->setOperand(0, adaptor.getSource()); });
return success();
}
};
diff --git a/compiler/src/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp b/compiler/src/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp
index e33d4c7..bbdc61f 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp
@@ -111,7 +111,7 @@
newResults.push_back(results[index]);
}
}
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
yieldOp, [&]() { yieldOp.getOperation()->setOperands(newResults); });
return newResults;
}
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/GPUTensorAlloc.cpp b/compiler/src/iree/compiler/Codegen/Common/GPU/GPUTensorAlloc.cpp
index 21c0123..c4a4acf 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/GPUTensorAlloc.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/GPUTensorAlloc.cpp
@@ -118,7 +118,7 @@
allocOp.getLoc(), allocOp.getType(), allocOp.getDynamicSizes(),
/*copy=*/Value(),
memorySpace ? cast<IntegerAttr>(*memorySpace) : IntegerAttr());
- rewriter.updateRootInPlace(linalgOp, [&]() {
+ rewriter.modifyOpInPlace(linalgOp, [&]() {
linalgOp->setOperand(linalgOp.getNumDpsInputs() + resultNumber,
newAllocOp);
});
diff --git a/compiler/src/iree/compiler/Codegen/Common/HoistUnrolledVectorExtractInsertSlice.cpp b/compiler/src/iree/compiler/Codegen/Common/HoistUnrolledVectorExtractInsertSlice.cpp
index 7fe3f76..6d5df32 100644
--- a/compiler/src/iree/compiler/Codegen/Common/HoistUnrolledVectorExtractInsertSlice.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/HoistUnrolledVectorExtractInsertSlice.cpp
@@ -126,10 +126,10 @@
if (!forOp.isDefinedOutsideOfLoop(extractStridedSliceOp.getVector())) {
assert(extractStridedSliceOp.getVector() == tensorBBArg &&
"extractSlice source not defined above must be the tracked bbArg");
- rewriter.startRootUpdate(extractStridedSliceOp);
+ rewriter.startOpModification(extractStridedSliceOp);
extractStridedSliceOp.getVectorMutable().assign(
forOp.getInitArgs()[initArgNumber]);
- rewriter.finalizeRootUpdate(extractStridedSliceOp);
+ rewriter.finalizeOpModification(extractStridedSliceOp);
}
}
@@ -149,20 +149,20 @@
// 3. Update the yield. Invariant: initArgNumber is the destination tensor.
auto yieldOp =
cast<scf::YieldOp>(newForOp.getRegion().front().getTerminator());
- rewriter.startRootUpdate(yieldOp);
+ rewriter.startOpModification(yieldOp);
yieldOp->setOperand(initArgNumber, insertOps[0].getDest());
- rewriter.finalizeRootUpdate(yieldOp);
+ rewriter.finalizeOpModification(yieldOp);
// 4. Hoist all the write ops after and make uses of
// newForOp.getResult(initArgNumber) flow through it.
for (auto [idx, insertStridedSliceOp] : llvm::enumerate(insertOps)) {
insertStridedSliceOp->moveAfter(newForOp);
- rewriter.startRootUpdate(insertStridedSliceOp);
+ rewriter.startOpModification(insertStridedSliceOp);
insertStridedSliceOp.getSourceMutable().assign(
newForOp.getResults()[initArgNumber + idx + 1]);
insertStridedSliceOp.getDestMutable().assign(
newForOp.getResults()[initArgNumber]);
- rewriter.finalizeRootUpdate(insertStridedSliceOp);
+ rewriter.finalizeOpModification(insertStridedSliceOp);
rewriter.replaceAllUsesExcept(newForOp.getResult(initArgNumber),
insertStridedSliceOp.getResult(),
insertStridedSliceOp);
diff --git a/compiler/src/iree/compiler/Codegen/Common/ReplaceSlowMinMaxOps.cpp b/compiler/src/iree/compiler/Codegen/Common/ReplaceSlowMinMaxOps.cpp
index 7ec3f27..1794388 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ReplaceSlowMinMaxOps.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ReplaceSlowMinMaxOps.cpp
@@ -40,13 +40,13 @@
LogicalResult matchAndRewrite(SlowReductionOp slowReductionOp,
PatternRewriter &rewriter) const override {
if (slowReductionOp.getKind() == vector::CombiningKind::MINIMUMF) {
- rewriter.updateRootInPlace(slowReductionOp, [&]() {
+ rewriter.modifyOpInPlace(slowReductionOp, [&]() {
slowReductionOp.setKind(vector::CombiningKind::MINNUMF);
});
return success();
}
if (slowReductionOp.getKind() == vector::CombiningKind::MAXIMUMF) {
- rewriter.updateRootInPlace(slowReductionOp, [&]() {
+ rewriter.modifyOpInPlace(slowReductionOp, [&]() {
slowReductionOp.setKind(vector::CombiningKind::MAXNUMF);
});
return success();
diff --git a/compiler/src/iree/compiler/Codegen/Common/TestPartitionableLoopsInterface.cpp b/compiler/src/iree/compiler/Codegen/Common/TestPartitionableLoopsInterface.cpp
index a45a202..749f17e 100644
--- a/compiler/src/iree/compiler/Codegen/Common/TestPartitionableLoopsInterface.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/TestPartitionableLoopsInterface.cpp
@@ -40,8 +40,8 @@
auto constantAttr = DenseIntElementsAttr::get(type, partitionableLoops);
rewriter.create<IREE::Util::UnfoldableConstantOp>(interfaceOp.getLoc(),
constantAttr);
- rewriter.updateRootInPlace(
- interfaceOp, [&] { interfaceOp->removeAttr(kAttributeName); });
+ rewriter.modifyOpInPlace(interfaceOp,
+ [&] { interfaceOp->removeAttr(kAttributeName); });
return success();
}
};
diff --git a/compiler/src/iree/compiler/Codegen/Common/TransformExtensions/CommonExtensions.cpp b/compiler/src/iree/compiler/Codegen/Common/TransformExtensions/CommonExtensions.cpp
index 008e033..ee0495c 100644
--- a/compiler/src/iree/compiler/Codegen/Common/TransformExtensions/CommonExtensions.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/TransformExtensions/CommonExtensions.cpp
@@ -380,7 +380,7 @@
}
// Promote extract_slice source to bbArg.
- rewriter.updateRootInPlace(extractSliceOp, [&]() {
+ rewriter.modifyOpInPlace(extractSliceOp, [&]() {
extractSliceOp.getSourceMutable().assign(bbArg);
});
}
@@ -475,7 +475,7 @@
// Step 6. RAUW thread indices to thread ops.
for (Value blockIdx : forallOp.getInductionVars()) {
for (Operation *user : llvm::make_early_inc_range(blockIdx.getUsers())) {
- rewriter.updateRootInPlace(user, [&]() {
+ rewriter.modifyOpInPlace(user, [&]() {
user->replaceUsesOfWith(blockIdx, bvm.lookup(blockIdx));
});
}
diff --git a/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp b/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp
index ceba477..05b9ae6 100644
--- a/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp
@@ -558,7 +558,7 @@
return rewriter.notifyMatchFailure(funcOp,
"failed to convert region types");
}
- rewriter.updateRootInPlace(funcOp, []() {});
+ rewriter.modifyOpInPlace(funcOp, []() {});
return success();
}
};
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUTileAndFuse.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUTileAndFuse.cpp
index a405cf5..949f4d7 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUTileAndFuse.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPUTileAndFuse.cpp
@@ -162,10 +162,10 @@
if (dpsOp.isDpsInit(&operand) &&
mapToIterArg.contains(sliceOp.getSource())) {
- rewriter.startRootUpdate(sliceOp);
+ rewriter.startOpModification(sliceOp);
sliceOp.getSourceMutable().assign(
mapToIterArg[sliceOp.getSource()]);
- rewriter.finalizeRootUpdate(sliceOp);
+ rewriter.finalizeOpModification(sliceOp);
}
}
};
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensions.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensions.cpp
index 063be1a..db54a2e 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensions.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensions.cpp
@@ -103,10 +103,10 @@
return diag;
auto newAttr = rewriter.getIndexArrayAttr(getWorkgroupDims());
auto subgroupSizeAttr = rewriter.getIndexAttr(getSubgroupSize());
- rewriter.startRootUpdate(exportOp);
+ rewriter.startOpModification(exportOp);
exportOp->setAttr(exportOp.getWorkgroupSizeAttrName(), newAttr);
exportOp->setAttr(exportOp.getSubgroupSizeAttrName(), subgroupSizeAttr);
- rewriter.finalizeRootUpdate(exportOp);
+ rewriter.finalizeOpModification(exportOp);
return DiagnosedSilenceableFailure::success();
}
@@ -148,9 +148,9 @@
for (Operation *user : llvm::make_early_inc_range(laneId.getUsers())) {
if (!executeOp->isProperAncestor(user))
continue;
- b.startRootUpdate(user);
+ b.startOpModification(user);
user->replaceUsesOfWith(laneId, zero);
- b.finalizeRootUpdate(user);
+ b.finalizeOpModification(user);
applied = true;
}
return success(applied);
@@ -480,14 +480,14 @@
// the yielded type depending on whether the op has "broadcast"
// behavior (see the doc of WarpExecuteOnLane0Op).
for (OpOperand &use : distributedVal.getUses()) {
- rewriter.startRootUpdate(use.getOwner());
+ rewriter.startOpModification(use.getOwner());
Value replacement = newRead;
if (use.get().getType() != newRead.getType()) {
replacement = rewriter.create<vector::BroadcastOp>(
load.getLoc(), use.get().getType(), newRead);
}
use.getOwner()->setOperand(use.getOperandNumber(), replacement);
- rewriter.finalizeRootUpdate(use.getOwner());
+ rewriter.finalizeOpModification(use.getOwner());
}
return success();
}
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp
index 59cf41d..d40bf4c 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp
@@ -323,7 +323,7 @@
signatureConverter);
// Creates a new function with the update signature.
- rewriter.updateRootInPlace(funcOp, [&] {
+ rewriter.modifyOpInPlace(funcOp, [&] {
funcOp.setType(rewriter.getFunctionType(
signatureConverter.getConvertedTypes(), std::nullopt));
});
@@ -735,8 +735,8 @@
LogicalResult
matchAndRewrite(OpT op, typename OpT::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- rewriter.updateRootInPlace(op,
- [&] { op->setOperands(adaptor.getOperands()); });
+ rewriter.modifyOpInPlace(op,
+ [&] { op->setOperands(adaptor.getOperands()); });
return success();
}
};
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp
index c318881..2b76307 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOpFolders.cpp
@@ -345,7 +345,7 @@
rewriter.getIndexAttr(
oldOrdinalPosToNewOrdinalPos.lookup(oldOrdinalPos)));
}
- rewriter.updateRootInPlace(op, []() {});
+ rewriter.modifyOpInPlace(op, []() {});
return success();
}
};
@@ -459,7 +459,7 @@
auto oldValues = llvm::to_vector(oldValueRange);
for (unsigned i = 0; i < dynamicDims.size(); ++i) {
if (oldValues[i] != dynamicDims[i]) {
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
op, [&]() { mutableDimValues.slice(i, 1).assign(dynamicDims[i]); });
anyChanged = true;
}
@@ -712,8 +712,8 @@
auto newAttr = deduplicateArrayElements(originalAttr);
if (newAttr == originalAttr)
return failure();
- rewriter.updateRootInPlace(
- dispatchOp, [&]() { dispatchOp.setEntryPointsAttr(newAttr); });
+ rewriter.modifyOpInPlace(dispatchOp,
+ [&]() { dispatchOp.setEntryPointsAttr(newAttr); });
return success();
}
};
diff --git a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.cpp b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.cpp
index 12e31b0..7d0b473 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/IR/FlowOps.cpp
@@ -623,7 +623,7 @@
for (const auto &it : llvm::enumerate(returnOp.getOperands()))
if (!unusedResults.contains(it.index()))
yieldedValues.push_back(it.value());
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
returnOp, [&]() { returnOp.getOperandsMutable().assign(yieldedValues); });
// Replace all uses of the old op.
diff --git a/compiler/src/iree/compiler/Dialect/Flow/TransformExtensions/FlowExtensions.cpp b/compiler/src/iree/compiler/Dialect/Flow/TransformExtensions/FlowExtensions.cpp
index 71cc462..b95d09e 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/TransformExtensions/FlowExtensions.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/TransformExtensions/FlowExtensions.cpp
@@ -215,7 +215,7 @@
uses.push_back(&use);
for (OpOperand *use : uses) {
unsigned resultNum = llvm::cast<OpResult>(use->get()).getResultNumber();
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
use->getOwner(), [&]() { use->set(cloned->getOpResult(resultNum)); });
}
}
@@ -728,7 +728,7 @@
if (operand.get().getDefiningOp() == regionOp) {
unsigned resultNumber =
llvm::cast<OpResult>(operand.get()).getResultNumber();
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
target, [&]() { operand.set(returnOp->getOperand(resultNumber)); });
}
}
@@ -749,7 +749,7 @@
// Replace uses of `target` after the dispatch region.
for (OpOperand *use : usesOutsideOfRegion) {
- rewriter.updateRootInPlace(use->getOwner(), [&]() {
+ rewriter.modifyOpInPlace(use->getOwner(), [&]() {
use->set(regionOp->getResult(
previousNumResults +
llvm::cast<OpResult>(use->get()).getResultNumber()));
diff --git a/compiler/src/iree/compiler/Dialect/Flow/Transforms/FormDispatchWorkgroups.cpp b/compiler/src/iree/compiler/Dialect/Flow/Transforms/FormDispatchWorkgroups.cpp
index 4eaa40d..0d6c694 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/Transforms/FormDispatchWorkgroups.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/Transforms/FormDispatchWorkgroups.cpp
@@ -225,7 +225,7 @@
rewriter.create<Flow::ReturnOp>(loc, defaultCountOp.getResults());
// Update the `workgroupsOp` region.
- rewriter.updateRootInPlace(workgroupsOp, [&]() {
+ rewriter.modifyOpInPlace(workgroupsOp, [&]() {
// Update the workload of the op.
workgroupsOp.getWorkloadMutable().assign(workload);
diff --git a/compiler/src/iree/compiler/Dialect/Flow/Transforms/RegionOpUtils.cpp b/compiler/src/iree/compiler/Dialect/Flow/Transforms/RegionOpUtils.cpp
index 54b9216..41389c8 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/Transforms/RegionOpUtils.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/Transforms/RegionOpUtils.cpp
@@ -391,7 +391,7 @@
// Replace all uses in the dispatch region.
for (OpOperand *use : usesInsideOfRegion) {
- rewriter.updateRootInPlace(use->getOwner(), [&]() {
+ rewriter.modifyOpInPlace(use->getOwner(), [&]() {
use->set(newTargetOp->getResult(
llvm::cast<OpResult>(use->get()).getResultNumber()));
});
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Conversion/StreamToHAL/Patterns.cpp b/compiler/src/iree/compiler/Dialect/HAL/Conversion/StreamToHAL/Patterns.cpp
index e4eadec..16f9e3e 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Conversion/StreamToHAL/Patterns.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Conversion/StreamToHAL/Patterns.cpp
@@ -1249,7 +1249,7 @@
return failure();
if (!llvm::isa<IREE::Stream::TimepointAttr>(*initialValue))
return failure();
- rewriter.updateRootInPlace(op, [&]() { op.removeInitialValueAttr(); });
+ rewriter.modifyOpInPlace(op, [&]() { op.removeInitialValueAttr(); });
return success();
}
};
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Conversion/UtilToHAL/Patterns.cpp b/compiler/src/iree/compiler/Dialect/HAL/Conversion/UtilToHAL/Patterns.cpp
index f947430..afc3349 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Conversion/UtilToHAL/Patterns.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/Conversion/UtilToHAL/Patterns.cpp
@@ -27,7 +27,7 @@
auto newType = getTypeConverter()->convertType(op.getType());
if (newType == op.getType())
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
// NOTE: the initial value may be invalid here! We rely on
// dialect-specific conversions to handle it.
op.setTypeAttr(TypeAttr::get(newType));
diff --git a/compiler/src/iree/compiler/Dialect/HAL/IR/HALOpFolders.cpp b/compiler/src/iree/compiler/Dialect/HAL/IR/HALOpFolders.cpp
index 349024d..371b6f8 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/IR/HALOpFolders.cpp
+++ b/compiler/src/iree/compiler/Dialect/HAL/IR/HALOpFolders.cpp
@@ -144,7 +144,7 @@
rewriter.restoreInsertionPoint(ip);
if (!needsUpdate)
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getSourceBufferMutable().assign(newSourceBuffer);
op.getSourceOffsetMutable().assign(newSourceOffset);
});
@@ -232,7 +232,7 @@
rewriter.restoreInsertionPoint(ip);
if (!needsUpdate)
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getTargetBufferMutable().assign(newTargetBuffer);
op.getTargetOffsetMutable().assign(newTargetOffset);
});
@@ -282,7 +282,7 @@
rewriter.restoreInsertionPoint(ip);
if (!needsUpdate)
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getSourceBufferMutable().assign(newSourceBuffer);
op.getSourceOffsetMutable().assign(newSourceOffset);
op.getTargetBufferMutable().assign(newTargetBuffer);
@@ -328,7 +328,7 @@
rewriter.restoreInsertionPoint(ip);
if (!needsUpdate)
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
auto mutableBindingBuffers = op.getBindingBuffersMutable();
mutableBindingBuffers.clear();
mutableBindingBuffers.append(bindingBuffers);
@@ -531,24 +531,24 @@
return rewriter.notifyMatchFailure(barrierOp, "signaling op not found");
}
- rewriter.startRootUpdate(signalingOp);
+ rewriter.startOpModification(signalingOp);
// Try to move the fence producer before the signaling op. This will fail if
// the op creating the fence has dependencies with hazards.
if (!IREE::Util::tryMoveProducerBefore(newFence, signalingOp)) {
- rewriter.cancelRootUpdate(signalingOp);
+ rewriter.cancelOpModification(signalingOp);
return rewriter.notifyMatchFailure(barrierOp,
"fence is not usable by signaling op");
}
// Rewrite the signaling op to signal the barrier fence.
if (failed(updateOpToSignalFence(signalingOp, newFence))) {
- rewriter.cancelRootUpdate(signalingOp);
+ rewriter.cancelOpModification(signalingOp);
return rewriter.notifyMatchFailure(barrierOp,
"unrecognized signaling op");
}
- rewriter.finalizeRootUpdate(signalingOp);
+ rewriter.finalizeOpModification(signalingOp);
// Elide the barrier. The fence should be cleaned up as part of DCE.
rewriter.eraseOp(barrierOp);
@@ -669,7 +669,7 @@
"not enough blocks to merge");
}
- rewriter.startRootUpdate(variantOp);
+ rewriter.startOpModification(variantOp);
// Gather all constants initialized by the blocks.
SmallVector<Location> blockLocs;
@@ -758,7 +758,7 @@
OpBuilder::atBlockBegin(postBlock).create<IREE::HAL::ReturnOp>(
fusedLoc, resultValues);
- rewriter.finalizeRootUpdate(variantOp);
+ rewriter.finalizeOpModification(variantOp);
// Erase all the old blocks.
for (auto blockOp : blockOps) {
@@ -802,7 +802,7 @@
auto deviceArg = blockOp.getArgument(0);
if (!deviceArg.use_empty())
return failure();
- rewriter.updateRootInPlace(blockOp, [&]() {
+ rewriter.modifyOpInPlace(blockOp, [&]() {
blockOp.eraseArgument(0);
blockOp.setFunctionTypeAttr(TypeAttr::get(
rewriter.getFunctionType(/*inputs=*/{}, blockOp.getResultTypes())));
@@ -839,7 +839,7 @@
}
// Update function in-place.
- rewriter.updateRootInPlace(blockOp, [&]() {
+ rewriter.modifyOpInPlace(blockOp, [&]() {
// Update metadata.
blockOp.setFunctionTypeAttr(TypeAttr::get(
rewriter.getFunctionType(blockOp.getArgumentTypes(), resultTypes)));
diff --git a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamDialect.cpp b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamDialect.cpp
index 15dfe32..2a47f03 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamDialect.cpp
+++ b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamDialect.cpp
@@ -80,8 +80,8 @@
dyn_cast<IREE::Stream::ResourceSizeOp>(use.getOwner())) {
rewriter.replaceOp(sizeOp, sizeValue);
} else {
- rewriter.updateRootInPlace(use.getOwner(),
- [&]() { use.set(resourceValue); });
+ rewriter.modifyOpInPlace(use.getOwner(),
+ [&]() { use.set(resourceValue); });
}
}
rewriter.eraseOp(castOp);
diff --git a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOpFolders.cpp b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOpFolders.cpp
index 810f547..7275e80 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOpFolders.cpp
+++ b/compiler/src/iree/compiler/Dialect/Stream/IR/StreamOpFolders.cpp
@@ -163,7 +163,7 @@
// Replace the pattern on the op with the new one.
auto narrowValue =
rewriter.create<arith::ConstantOp>(fillOp.getLoc(), newPatternAttr);
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
fillOp, [&]() { fillOp.getValueMutable().assign(narrowValue); });
return success();
}
@@ -475,7 +475,7 @@
IREE::Util::TiedOpInterface::findTiedBaseValue(result.value());
if (auto blockArg = llvm::dyn_cast<BlockArgument>(baseValue)) {
unsigned operandIndex = blockArg.getArgNumber();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.setTiedResultOperandIndex(result.index(), operandIndex);
});
didModify = true;
@@ -527,8 +527,8 @@
op.getAwaitTimepoint().getDefiningOp());
if (!isImmediate)
return failure();
- rewriter.updateRootInPlace(
- op, [&]() { op.getAwaitTimepointMutable().clear(); });
+ rewriter.modifyOpInPlace(op,
+ [&]() { op.getAwaitTimepointMutable().clear(); });
return success();
}
};
@@ -560,7 +560,7 @@
}
if (replacements.empty())
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
auto newTimepoint = joinAwaitTimepoints(
op.getLoc(), op.getAwaitTimepoint(), newTimepoints, rewriter);
op.getAwaitTimepointMutable().assign(newTimepoint);
@@ -690,7 +690,7 @@
auto fusedLoc = rewriter.getFusedLoc({subviewOp.getLoc(), op.getLoc()});
auto newOffset = rewriter.createOrFold<arith::AddIOp>(
fusedLoc, subviewOp.getSourceOffset(), op.getSourceOffset());
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getSourceMutable().assign(subviewOp.getSource());
op.getSourceSizeMutable().assign(subviewOp.getSourceSize());
op.getSourceOffsetMutable().assign(newOffset);
@@ -735,7 +735,7 @@
auto fusedLoc = rewriter.getFusedLoc({subviewOp.getLoc(), op.getLoc()});
auto newOffset = rewriter.createOrFold<arith::AddIOp>(
fusedLoc, subviewOp.getSourceOffset(), op.getTargetOffset());
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getTargetMutable().assign(subviewOp.getSource());
op.getTargetSizeMutable().assign(subviewOp.getSourceSize());
op.getTargetOffsetMutable().assign(newOffset);
@@ -801,7 +801,7 @@
return failure();
// We always strip the offset here.
- rewriter.updateRootInPlace(op, [&]() { op.getOffsetMutable().clear(); });
+ rewriter.modifyOpInPlace(op, [&]() { op.getOffsetMutable().clear(); });
// Zero offsets don't do anything and can just be removed so we can avoid
// inserting a bunch of additional IR.
@@ -1024,7 +1024,7 @@
rewriter.restoreInsertionPoint(ip);
if (!needsUpdate)
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getSourceOffsetsMutable().assign(newSourceOffsets);
op.getResultSizesMutable().assign(newResultSizes);
});
@@ -1075,7 +1075,7 @@
rewriter.restoreInsertionPoint(ip);
if (!needsUpdate)
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getSourceOffsetMutable().assign(newSourceOffset);
op.getTargetMutable().assign(newTargetResource);
op.getTargetSizeMutable().assign(newTargetSize);
@@ -1128,7 +1128,7 @@
rewriter.restoreInsertionPoint(ip);
if (!needsUpdate)
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getSourceMutable().assign(newSourceResource);
op.getSourceSizeMutable().assign(newSourceSize);
op.getSourceOffsetMutable().assign(newSourceOffset);
@@ -1762,7 +1762,7 @@
fusedLoc, fillOp.getTargetLength(), sourceOp.getTargetLength());
}
- rewriter.updateRootInPlace(fillOp, [&]() {
+ rewriter.modifyOpInPlace(fillOp, [&]() {
sourceOp.getTargetOffsetMutable().assign(newOffset);
sourceOp.getTargetEndMutable().assign(newEnd);
sourceOp.getTargetLengthMutable().assign(newLength);
@@ -2012,7 +2012,7 @@
dyn_cast<arith::BitcastOp>(*loadedValue.getUsers().begin());
if (!bitcastOp)
return failure();
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
loadOp, [&]() { loadedValue.setType(bitcastOp.getType()); });
rewriter.replaceOp(bitcastOp, loadedValue);
return success();
@@ -2046,7 +2046,7 @@
auto storedValue = storeOp.getValue();
if (auto bitcastOp =
dyn_cast_or_null<arith::BitcastOp>(storedValue.getDefiningOp())) {
- rewriter.updateRootInPlace(storeOp, [&]() {
+ rewriter.modifyOpInPlace(storeOp, [&]() {
storeOp.getValueMutable().assign(bitcastOp.getOperand());
});
return success();
@@ -2079,8 +2079,8 @@
auto newAttr = deduplicateArrayElements(originalAttr);
if (newAttr == originalAttr)
return failure();
- rewriter.updateRootInPlace(
- dispatchOp, [&]() { dispatchOp.setEntryPointsAttr(newAttr); });
+ rewriter.modifyOpInPlace(dispatchOp,
+ [&]() { dispatchOp.setEntryPointsAttr(newAttr); });
return success();
}
};
@@ -2132,7 +2132,7 @@
}
if (captures.empty())
return failure();
- rewriter.startRootUpdate(op);
+ rewriter.startOpModification(op);
auto &entryBlock = op.getBody().front();
rewriter.setInsertionPointToStart(&entryBlock);
@@ -2155,7 +2155,7 @@
rewriter.replaceAllUsesExcept(arg, newOp.getResult(), newOp);
}
- rewriter.finalizeRootUpdate(op);
+ rewriter.finalizeOpModification(op);
return success();
}
};
@@ -2247,7 +2247,7 @@
auto fusedLoc = rewriter.getFusedLoc({subviewOp.getLoc(), op.getLoc()});
auto newOffset = rewriter.createOrFold<arith::AddIOp>(
fusedLoc, subviewOp.getSourceOffset(), op.getTargetOffset());
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getTargetMutable().assign(subviewOp.getSource());
op.getTargetSizeMutable().assign(subviewOp.getSourceSize());
op.getTargetOffsetMutable().assign(newOffset);
@@ -2289,7 +2289,7 @@
auto fusedLoc = rewriter.getFusedLoc({subviewOp.getLoc(), op.getLoc()});
auto newOffset = rewriter.createOrFold<arith::AddIOp>(
fusedLoc, subviewOp.getSourceOffset(), op.getTargetOffset());
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getTargetMutable().assign(subviewOp.getSource());
op.getTargetSizeMutable().assign(subviewOp.getSourceSize());
op.getTargetOffsetMutable().assign(newOffset);
@@ -2330,7 +2330,7 @@
auto fusedLoc = rewriter.getFusedLoc({subviewOp.getLoc(), op.getLoc()});
auto newOffset = rewriter.createOrFold<arith::AddIOp>(
fusedLoc, subviewOp.getSourceOffset(), op.getTargetOffset());
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getTargetMutable().assign(subviewOp.getSource());
op.getTargetSizeMutable().assign(subviewOp.getSourceSize());
op.getTargetOffsetMutable().assign(newOffset);
@@ -2371,7 +2371,7 @@
auto fusedLoc = rewriter.getFusedLoc({subviewOp.getLoc(), op.getLoc()});
auto newOffset = rewriter.createOrFold<arith::AddIOp>(
fusedLoc, subviewOp.getSourceOffset(), op.getTargetOffset());
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getTargetMutable().assign(subviewOp.getSource());
op.getTargetSizeMutable().assign(subviewOp.getSourceSize());
op.getTargetOffsetMutable().assign(newOffset);
@@ -2416,7 +2416,7 @@
rewriter.getFusedLoc({sourceSubviewOp.getLoc(), op.getLoc()});
auto newOffset = rewriter.createOrFold<arith::AddIOp>(
fusedLoc, sourceSubviewOp.getSourceOffset(), op.getSourceOffset());
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getSourceMutable().assign(sourceSubviewOp.getSource());
op.getSourceSizeMutable().assign(sourceSubviewOp.getSourceSize());
op.getSourceOffsetMutable().assign(newOffset);
@@ -2427,7 +2427,7 @@
rewriter.getFusedLoc({targetSubviewOp.getLoc(), op.getLoc()});
auto newOffset = rewriter.createOrFold<arith::AddIOp>(
fusedLoc, targetSubviewOp.getSourceOffset(), op.getTargetOffset());
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getTargetMutable().assign(targetSubviewOp.getSource());
op.getTargetSizeMutable().assign(targetSubviewOp.getSourceSize());
op.getTargetOffsetMutable().assign(newOffset);
@@ -2480,7 +2480,7 @@
}
if (!anySubviewOps)
return failure();
- rewriter.startRootUpdate(op);
+ rewriter.startOpModification(op);
setInsertionPointToParentExecutionScope(op, rewriter);
for (auto [resourceIndex, subviewOp] :
@@ -2500,7 +2500,7 @@
op.getResourceOffsetsMutable().slice(resourceIndex, 1).assign(newOffset);
}
- rewriter.finalizeRootUpdate(op);
+ rewriter.finalizeOpModification(op);
return success();
}
};
@@ -2527,8 +2527,8 @@
auto newAttr = deduplicateArrayElements(originalAttr);
if (newAttr == originalAttr)
return failure();
- rewriter.updateRootInPlace(
- dispatchOp, [&]() { dispatchOp.setEntryPointsAttr(newAttr); });
+ rewriter.modifyOpInPlace(dispatchOp,
+ [&]() { dispatchOp.setEntryPointsAttr(newAttr); });
return success();
}
};
@@ -2568,7 +2568,7 @@
}
if (!anySubviewOps)
return failure();
- rewriter.startRootUpdate(op);
+ rewriter.startOpModification(op);
setInsertionPointToParentExecutionScope(op, rewriter);
for (auto [resourceIndex, resourceSubviewOp] :
@@ -2591,7 +2591,7 @@
.assign(newOffset);
}
- rewriter.finalizeRootUpdate(op);
+ rewriter.finalizeOpModification(op);
return success();
}
};
@@ -2639,7 +2639,7 @@
}
if (captures.empty())
return failure();
- rewriter.startRootUpdate(op);
+ rewriter.startOpModification(op);
auto &entryBlock = op.getBody().front();
rewriter.setInsertionPointToStart(&entryBlock);
@@ -2662,7 +2662,7 @@
rewriter.replaceAllUsesExcept(arg, newOp.getResult(), newOp);
}
- rewriter.finalizeRootUpdate(op);
+ rewriter.finalizeOpModification(op);
return success();
}
};
@@ -2881,7 +2881,7 @@
rewriter.replaceOpWithNewOp<TimepointImmediateOp>(
op, op.getResultTimepoint().getType());
} else {
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
op, [&]() { op.getAwaitTimepointsMutable().assign(newTimepoints); });
}
return success();
@@ -2898,7 +2898,7 @@
op.getAwaitTimepoints().end());
if (newTimepoints.size() == op.getAwaitTimepoints().size())
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getAwaitTimepointsMutable().assign(newTimepoints.takeVector());
});
return success();
@@ -2929,7 +2929,7 @@
}
if (!didExpand)
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getAwaitTimepointsMutable().assign(newTimepoints.takeVector());
});
return success();
@@ -3131,8 +3131,8 @@
if (!canStablySinkTo(op, firstUserInDominator))
return failure();
- rewriter.updateRootInPlace(op,
- [&]() { op->moveBefore(firstUserInDominator); });
+ rewriter.modifyOpInPlace(op,
+ [&]() { op->moveBefore(firstUserInDominator); });
return success();
}
};
@@ -3144,7 +3144,7 @@
using OpRewritePattern::OpRewritePattern;
LogicalResult matchAndRewrite(TimepointAwaitOp op,
PatternRewriter &rewriter) const override {
- rewriter.startRootUpdate(op);
+ rewriter.startOpModification(op);
bool didChange = false;
for (auto operand : llvm::enumerate(op.getResourceOperands())) {
auto subviewOp =
@@ -3173,10 +3173,10 @@
.assign(subviewOp.getSource());
}
if (didChange) {
- rewriter.finalizeRootUpdate(op);
+ rewriter.finalizeOpModification(op);
return success();
} else {
- rewriter.cancelRootUpdate(op);
+ rewriter.cancelOpModification(op);
return failure();
}
}
diff --git a/compiler/src/iree/compiler/Dialect/Stream/Transforms/ConvertToStream.cpp b/compiler/src/iree/compiler/Dialect/Stream/Transforms/ConvertToStream.cpp
index 422e0ca..7350487 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/Transforms/ConvertToStream.cpp
+++ b/compiler/src/iree/compiler/Dialect/Stream/Transforms/ConvertToStream.cpp
@@ -161,7 +161,7 @@
newOperands.push_back(buildTensorExportOp(
op->getLoc(), newOperand, tensorType, dynamicDims, rewriter));
}
- rewriter.updateRootInPlace(op, [&]() { op->setOperands(newOperands); });
+ rewriter.modifyOpInPlace(op, [&]() { op->setOperands(newOperands); });
// Import into resources from tensor results produced by the op.
rewriter.setInsertionPointAfter(op);
diff --git a/compiler/src/iree/compiler/Dialect/Stream/Transforms/EncodeTensors.cpp b/compiler/src/iree/compiler/Dialect/Stream/Transforms/EncodeTensors.cpp
index e714b44..94c0f6e 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/Transforms/EncodeTensors.cpp
+++ b/compiler/src/iree/compiler/Dialect/Stream/Transforms/EncodeTensors.cpp
@@ -653,8 +653,8 @@
// Directly swap the type with the one, changing all uses in the IR.
// This works because
- rewriter.updateRootInPlace(op,
- [&]() { op.getResult().setType(alignedType); });
+ rewriter.modifyOpInPlace(op,
+ [&]() { op.getResult().setType(alignedType); });
return success();
}
@@ -686,7 +686,7 @@
rewriter.setInsertionPointAfterValue(loadedValue);
auto truncOp =
rewriter.create<arith::TruncIOp>(op.getLoc(), targetType, loadedValue);
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
loadedValue.replaceAllUsesExcept(truncOp, truncOp);
loadedValue.setType(alignedType);
});
@@ -718,7 +718,7 @@
// Extend the sub-byte -> byte type; e.g. i1 -> i8.
auto extOp = rewriter.create<arith::ExtUIOp>(op.getLoc(), alignedType,
op.getValue());
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
op, [&]() { op.getValueMutable().assign(extOp.getResult()); });
return success();
}
diff --git a/compiler/src/iree/compiler/Dialect/Stream/Transforms/RefineUsage.cpp b/compiler/src/iree/compiler/Dialect/Stream/Transforms/RefineUsage.cpp
index bf882a7..abb4cce 100644
--- a/compiler/src/iree/compiler/Dialect/Stream/Transforms/RefineUsage.cpp
+++ b/compiler/src/iree/compiler/Dialect/Stream/Transforms/RefineUsage.cpp
@@ -188,7 +188,7 @@
// Returns true if a change was made.
bool applyRegionTransitions(Operation *op, PatternRewriter &rewriter) const {
bool didChange = false;
- rewriter.startRootUpdate(op);
+ rewriter.startOpModification(op);
for (auto ®ion : op->getRegions()) {
for (auto &block : region) {
rewriter.setInsertionPoint(&block, block.begin());
@@ -200,9 +200,9 @@
}
}
if (didChange) {
- rewriter.finalizeRootUpdate(op);
+ rewriter.finalizeOpModification(op);
} else {
- rewriter.cancelRootUpdate(op);
+ rewriter.cancelOpModification(op);
}
return didChange;
}
@@ -345,7 +345,7 @@
LogicalResult matchAndRewrite(Op op,
PatternRewriter &rewriter) const override {
bool didChange = this->applyRegionTransitions(op, rewriter);
- rewriter.startRootUpdate(op);
+ rewriter.startOpModification(op);
rewriter.setInsertionPointAfter(op);
for (unsigned i = 0; i < op->getNumResults(); ++i) {
auto result = op->getResult(i);
@@ -355,9 +355,9 @@
}
}
if (didChange) {
- rewriter.finalizeRootUpdate(op);
+ rewriter.finalizeOpModification(op);
} else {
- rewriter.cancelRootUpdate(op);
+ rewriter.cancelOpModification(op);
}
return success(didChange);
}
@@ -376,7 +376,7 @@
bool didChange = this->applyRegionTransitions(op, rewriter);
auto affinityAttr = getOpAffinity(op);
- rewriter.startRootUpdate(op);
+ rewriter.startOpModification(op);
rewriter.setInsertionPointAfter(op);
auto sizeAwareOp =
@@ -394,9 +394,9 @@
}
if (didChange) {
- rewriter.finalizeRootUpdate(op);
+ rewriter.finalizeOpModification(op);
} else {
- rewriter.cancelRootUpdate(op);
+ rewriter.cancelOpModification(op);
}
return success(didChange);
}
diff --git a/compiler/src/iree/compiler/Dialect/Util/Conversion/ConversionPatterns.cpp b/compiler/src/iree/compiler/Dialect/Util/Conversion/ConversionPatterns.cpp
index 20373a6..041e6da 100644
--- a/compiler/src/iree/compiler/Dialect/Util/Conversion/ConversionPatterns.cpp
+++ b/compiler/src/iree/compiler/Dialect/Util/Conversion/ConversionPatterns.cpp
@@ -82,14 +82,14 @@
matchAndRewrite(IREE::Util::InitializerOp initializerOp, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
auto &typeConverter = *getTypeConverter();
- rewriter.startRootUpdate(initializerOp);
+ rewriter.startOpModification(initializerOp);
if (failed(rewriter.convertRegionTypes(&initializerOp.getBody(),
typeConverter))) {
- rewriter.cancelRootUpdate(initializerOp);
+ rewriter.cancelOpModification(initializerOp);
return rewriter.notifyMatchFailure(initializerOp,
"failed to convert region types");
}
- rewriter.finalizeRootUpdate(initializerOp);
+ rewriter.finalizeOpModification(initializerOp);
return success();
}
};
diff --git a/compiler/src/iree/compiler/Dialect/Util/IR/UtilOpFolders.cpp b/compiler/src/iree/compiler/Dialect/Util/IR/UtilOpFolders.cpp
index f76c1ff..b53de8f 100644
--- a/compiler/src/iree/compiler/Dialect/Util/IR/UtilOpFolders.cpp
+++ b/compiler/src/iree/compiler/Dialect/Util/IR/UtilOpFolders.cpp
@@ -565,7 +565,7 @@
auto globalOp =
SymbolTable::lookupNearestSymbolFrom<IREE::Util::GlobalOpInterface>(
storeOp->getParentOp(), storeOp.getGlobalAttr());
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
globalOp, [&]() { globalOp.setGlobalInitialValue(valueAttr); });
deadOps.push_back(storeOp);
@@ -744,7 +744,7 @@
fusedLoc, op.getSourceOffset(), oldRange.offset);
auto newRange = SubrangeOperand{op.getSource(), op.getSourceSize(),
newOffset, oldRange.length};
- rewriter.updateRootInPlace(subrangeOp, [&]() {
+ rewriter.modifyOpInPlace(subrangeOp, [&]() {
subrangeOp.setSubrangeOperand(use.getOperandNumber(), newRange);
});
}
@@ -893,7 +893,7 @@
rewriter.setInsertionPointAfter(op);
auto newOffset = rewriter.createOrFold<arith::AddIOp>(
fusedLoc, subspanOp.getSourceOffset(), op.getOffset());
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getOperandMutable().assign(subspanOp.getSource());
op.getOperandSizeMutable().assign(subspanOp.getSourceSize());
SmallPtrSet<Operation *, 2> exceptions;
diff --git a/compiler/src/iree/compiler/Dialect/Util/Transforms/Patterns.cpp b/compiler/src/iree/compiler/Dialect/Util/Transforms/Patterns.cpp
index 0ea3f10..b11761b 100644
--- a/compiler/src/iree/compiler/Dialect/Util/Transforms/Patterns.cpp
+++ b/compiler/src/iree/compiler/Dialect/Util/Transforms/Patterns.cpp
@@ -89,7 +89,7 @@
return failure(); // no dupes at all
}
- rewriter.startRootUpdate(op);
+ rewriter.startOpModification(op);
// Iterate over all blocks after the entry block. We can't change the entry
// block as it is part of the function signature.
@@ -154,7 +154,7 @@
auto operands = successorOperands.slice(
successorOperands.getProducedOperandCount(),
successorOperands.size());
- rewriter.updateRootInPlace(blockSource.branchOp, [&]() {
+ rewriter.modifyOpInPlace(blockSource.branchOp, [&]() {
eraseOperands(operands, elidedArgs);
});
}
@@ -164,10 +164,10 @@
}
if (didChange) {
- rewriter.finalizeRootUpdate(op);
+ rewriter.finalizeOpModification(op);
return success();
} else {
- rewriter.cancelRootUpdate(op);
+ rewriter.cancelOpModification(op);
return failure();
}
}
@@ -216,7 +216,7 @@
}
}
- rewriter.startRootUpdate(op);
+ rewriter.startOpModification(op);
// Iterate over all blocks after the entry block. We can't change the entry
// block as it is part of the function signature.
@@ -290,7 +290,7 @@
auto operands =
successorOperands.slice(successorOperands.getProducedOperandCount(),
successorOperands.size());
- rewriter.updateRootInPlace(blockSource.branchOp, [&]() {
+ rewriter.modifyOpInPlace(blockSource.branchOp, [&]() {
eraseOperands(operands, elidedArgs);
});
}
@@ -299,10 +299,10 @@
}
if (didChange) {
- rewriter.finalizeRootUpdate(op);
+ rewriter.finalizeOpModification(op);
return success();
} else {
- rewriter.cancelRootUpdate(op);
+ rewriter.cancelOpModification(op);
return failure();
}
}
diff --git a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp
index 31fa545..ca276f5 100644
--- a/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp
+++ b/compiler/src/iree/compiler/Dialect/VM/Conversion/VMToEmitC/ConvertVMToEmitC.cpp
@@ -1298,7 +1298,7 @@
signatureConverter);
// Creates a new function with the updated signature.
- rewriter.updateRootInPlace(funcOp, [&] {
+ rewriter.modifyOpInPlace(funcOp, [&] {
funcOp.setType(
rewriter.getFunctionType(signatureConverter.getConvertedTypes(),
funcOp.getFunctionType().getResults()));
diff --git a/compiler/src/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp b/compiler/src/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp
index 2905d15..0d36737 100644
--- a/compiler/src/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp
+++ b/compiler/src/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp
@@ -96,7 +96,7 @@
auto globalOp =
SymbolTable::lookupNearestSymbolFrom<IREE::Util::GlobalOpInterface>(
op, globalRefAttr);
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
globalOp, [&]() { globalOp.setGlobalInitialValue(valueAttr); });
deadOps.push_back(op);
});
diff --git a/compiler/src/iree/compiler/GlobalOptimization/DetachElementwiseFromNamedOps.cpp b/compiler/src/iree/compiler/GlobalOptimization/DetachElementwiseFromNamedOps.cpp
index 49b6c48..1da073c 100644
--- a/compiler/src/iree/compiler/GlobalOptimization/DetachElementwiseFromNamedOps.cpp
+++ b/compiler/src/iree/compiler/GlobalOptimization/DetachElementwiseFromNamedOps.cpp
@@ -79,8 +79,8 @@
rewriter.create<linalg::FillOp>(loc, zero, initOp.getResult()).result();
// Update the contraction op to use the new zero tensor as output operand.
- rewriter.updateRootInPlace(linalgOp,
- [&]() { linalgOp.setDpsInitOperand(0, fill); });
+ rewriter.modifyOpInPlace(linalgOp,
+ [&]() { linalgOp.setDpsInitOperand(0, fill); });
auto outputMap = mlir::compressUnusedDims(
linalgOp.getMatchingIndexingMap(outputOperands.front()));
@@ -175,7 +175,7 @@
.create<linalg::FillOp>(
loc, resultType, scalarConstantOp, emptyTensorOp)
.getResult(0);
- rewriter.updateRootInPlace(dpsInterfaceOp, [&]() {
+ rewriter.modifyOpInPlace(dpsInterfaceOp, [&]() {
dpsInterfaceOp.setDpsInitOperand(outOperand.index(), fillOp);
});
madeChanges = true;
diff --git a/compiler/src/iree/compiler/GlobalOptimization/PropagateLinalgTranspose.cpp b/compiler/src/iree/compiler/GlobalOptimization/PropagateLinalgTranspose.cpp
index 741992b..29ec576 100644
--- a/compiler/src/iree/compiler/GlobalOptimization/PropagateLinalgTranspose.cpp
+++ b/compiler/src/iree/compiler/GlobalOptimization/PropagateLinalgTranspose.cpp
@@ -352,7 +352,7 @@
// To do the fusion, we can simply apply the permutation of the transpose
// to the results of the associated input's indexing map, and then forward
// the input to the transpose to the consumer generic.
- rewriter.startRootUpdate(genericOp);
+ rewriter.startOpModification(genericOp);
SmallVector<AffineMap> newIndexingMaps = genericOp.getIndexingMapsArray();
AffineMap inputMap = genericOp.getMatchingIndexingMap(transposeOperand);
@@ -366,7 +366,7 @@
rewriter.getAffineMapArrayAttr(newIndexingMaps));
genericOp.setOperand(inputIndex, transposeOp.getInput());
- rewriter.finalizeRootUpdate(genericOp);
+ rewriter.finalizeOpModification(genericOp);
return success();
}
};
diff --git a/compiler/src/iree/compiler/GlobalOptimization/RemoveZeroExtentTensors.cpp b/compiler/src/iree/compiler/GlobalOptimization/RemoveZeroExtentTensors.cpp
index c6ee5c4..9cb5929 100644
--- a/compiler/src/iree/compiler/GlobalOptimization/RemoveZeroExtentTensors.cpp
+++ b/compiler/src/iree/compiler/GlobalOptimization/RemoveZeroExtentTensors.cpp
@@ -51,7 +51,7 @@
auto shape = tensor::getMixedSizes(rewriter, loc, operand.get());
auto emptyTensorOp = rewriter.create<tensor::EmptyOp>(
loc, shape, operandType->getElementType());
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
owner, [&]() { owner->setOperand(operandNum, emptyTensorOp); });
didUpdate = true;
}
diff --git a/compiler/src/iree/compiler/Modules/HAL/Inline/Conversion/StreamToHALInline/Patterns.cpp b/compiler/src/iree/compiler/Modules/HAL/Inline/Conversion/StreamToHALInline/Patterns.cpp
index a917de7..12c1ba0 100644
--- a/compiler/src/iree/compiler/Modules/HAL/Inline/Conversion/StreamToHALInline/Patterns.cpp
+++ b/compiler/src/iree/compiler/Modules/HAL/Inline/Conversion/StreamToHALInline/Patterns.cpp
@@ -624,7 +624,7 @@
return failure();
if (!llvm::isa<IREE::Stream::TimepointAttr>(*initialValue))
return failure();
- rewriter.updateRootInPlace(
+ rewriter.modifyOpInPlace(
op, [&]() { op.setInitialValueAttr(rewriter.getI64IntegerAttr(0)); });
return success();
}
diff --git a/compiler/src/iree/compiler/Modules/HAL/Inline/IR/HALInlineOps.cpp b/compiler/src/iree/compiler/Modules/HAL/Inline/IR/HALInlineOps.cpp
index d1a6802..74eb5fd 100644
--- a/compiler/src/iree/compiler/Modules/HAL/Inline/IR/HALInlineOps.cpp
+++ b/compiler/src/iree/compiler/Modules/HAL/Inline/IR/HALInlineOps.cpp
@@ -171,7 +171,7 @@
rewriter.restoreInsertionPoint(ip);
if (!needsUpdate)
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getSourceBufferMutable().assign(newSourceBuffer);
op.getSourceOffsetMutable().assign(newSourceOffset);
});
diff --git a/compiler/src/iree/compiler/Modules/HAL/Loader/IR/HALLoaderOps.cpp b/compiler/src/iree/compiler/Modules/HAL/Loader/IR/HALLoaderOps.cpp
index a355fea..467a339 100644
--- a/compiler/src/iree/compiler/Modules/HAL/Loader/IR/HALLoaderOps.cpp
+++ b/compiler/src/iree/compiler/Modules/HAL/Loader/IR/HALLoaderOps.cpp
@@ -181,7 +181,7 @@
}
if (!didChangeAny)
return failure();
- rewriter.updateRootInPlace(op, [&]() {
+ rewriter.modifyOpInPlace(op, [&]() {
op.getBindingBuffersMutable().assign(bindingBuffers);
op.getBindingOffsetsMutable().assign(bindingOffsets);
});
diff --git a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Transforms/Transforms.h b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Transforms/Transforms.h
index f2ed269..bdd7103 100644
--- a/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Transforms/Transforms.h
+++ b/llvm-external-projects/iree-dialects/include/iree-dialects/Dialect/LinalgExt/Transforms/Transforms.h
@@ -106,14 +106,14 @@
// So to fail properly, we should be cloning
// the op and deleting the previous op. This
// needs more investigation.
- rewriter.startRootUpdate(op);
+ rewriter.startOpModification(op);
std::optional<linalg::LinalgOp> promotedOp =
promoteSubViews(rewriter, cast<linalg::LinalgOp>(op), options);
if (!promotedOp) {
- rewriter.cancelRootUpdate(op);
+ rewriter.cancelOpModification(op);
return op->emitError("subview promotion failed");
}
- rewriter.finalizeRootUpdate(op);
+ rewriter.finalizeOpModification(op);
filter.replaceLinalgTransformationFilter(rewriter, op);
return success();
}
diff --git a/third_party/llvm-project b/third_party/llvm-project
index 905caa5..46537e6 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit 905caa57a4bf51db81a414056df5fba47adf2b21
+Subproject commit 46537e6a745a59b4e11ef87a5d494fe9075144ab
diff --git a/third_party/torch-mlir b/third_party/torch-mlir
index 6096fcb..65f75ab 160000
--- a/third_party/torch-mlir
+++ b/third_party/torch-mlir
@@ -1 +1 @@
-Subproject commit 6096fcb347691982d721d74d96794ac0d17af0d9
+Subproject commit 65f75ab7124fa000433aab2709553ec91a22fdde
