LLVM integrate to ebb78a95cede526ece9b904e9ba623d4b963df60 + cherry-pick (#9729)
* LLVM integrate
MLIR-MHLO: 7365a774d0c772e5a032ac0d36a1d532f17860f3
LLVM: ebb78a95cede526ece9b904e9ba623d4b963df60
Additionally cherry-picks the following IREE-local revert by dcaballe@:
* This reverts commit 1178992c72b002c3b2c87203252c566eeb273cc1
* Update scf::ParallelInsertSliceOp -> tensor::ParallelInsertSliceOp
* Update isLegalNumpyRankedBroadcast -> IsLegalNumpyRankedBroadcast
* Update dynamic-update-slice -> dynamic_update_slice
* Patch tests related to #hal.descriptor_set.layout / #hal.descriptor_set.binding
diff --git a/compiler/src/iree/compiler/Dialect/Flow/TransformExtensions/FlowExtensions.cpp b/compiler/src/iree/compiler/Dialect/Flow/TransformExtensions/FlowExtensions.cpp
index 6b477d3..18eff71 100644
--- a/compiler/src/iree/compiler/Dialect/Flow/TransformExtensions/FlowExtensions.cpp
+++ b/compiler/src/iree/compiler/Dialect/Flow/TransformExtensions/FlowExtensions.cpp
@@ -86,8 +86,8 @@
Location loc = performConcurrentlyOp.getLoc();
int64_t resultIndex = 0;
for (const Operation &yieldingOp :
- llvm::make_early_inc_range(performConcurrentlyOp.yieldingOps())) {
- auto parallelInsertOp = cast<scf::ParallelInsertSliceOp>(&yieldingOp);
+ llvm::make_early_inc_range(performConcurrentlyOp.getYieldingOps())) {
+ auto parallelInsertOp = cast<tensor::ParallelInsertSliceOp>(&yieldingOp);
OpBuilder::InsertionGuard g(rewriter);
rewriter.setInsertionPoint(block.getTerminator());
auto dynamicDims = Util::findVariadicDynamicDims(
@@ -179,8 +179,8 @@
// over to the Flow::DispatchWorkgroupsOp.
// Use a SetVector to ensure tensor operand uniqueness.
llvm::SetVector<Value> resultTensorOperands, resultTensorsDynamicDims;
- for (const Operation &yieldingOp : performConcurrentlyOp.yieldingOps()) {
- auto parallelInsertOp = cast<scf::ParallelInsertSliceOp>(&yieldingOp);
+ for (const Operation &yieldingOp : performConcurrentlyOp.getYieldingOps()) {
+ auto parallelInsertOp = cast<tensor::ParallelInsertSliceOp>(&yieldingOp);
Value dest = parallelInsertOp.getDest();
bool inserted = resultTensorOperands.insert(dest);
if (!inserted) continue;
diff --git a/compiler/src/iree/compiler/Dialect/HAL/IR/test/attributes.mlir b/compiler/src/iree/compiler/Dialect/HAL/IR/test/attributes.mlir
index 78e0a7b..7502616 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/IR/test/attributes.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/IR/test/attributes.mlir
@@ -13,11 +13,11 @@
// CHECK-LABEL: executable_layout.basic
"executable_layout.basic"() {
// CHECK: layout0 = #hal.executable.layout<push_constants = 4, sets = [
- // CHECK-SAME: #hal.descriptor_set.layout<0, bindings = [
- // CHECK-SAME: #hal.descriptor_set.binding<0, storage_buffer>
- // CHECK-SAME: #hal.descriptor_set.binding<1, storage_buffer>
- // CHECK-SAME: #hal.descriptor_set.layout<1, bindings = [
- // CHECK-SAME: #hal.descriptor_set.binding<0, uniform_buffer>
+ // CHECK-SAME: <0, bindings = [
+ // CHECK-SAME: <0, storage_buffer>
+ // CHECK-SAME: <1, storage_buffer>
+ // CHECK-SAME: <1, bindings = [
+ // CHECK-SAME: <0, uniform_buffer>
layout0 = #hal.executable.layout<push_constants = 4, sets = [
#hal.descriptor_set.layout<0, bindings = [
#hal.descriptor_set.binding<0, storage_buffer>,
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir b/compiler/src/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir
index c0e93e7..8770050 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir
+++ b/compiler/src/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir
@@ -41,10 +41,10 @@
// CHECK-NEXT: hal.executable.variant public @vmvx_bytecode_fb, target = <"vmvx", "vmvx-bytecode-fb"> {
// CHECK-NEXT: hal.executable.export public @add_dispatch_0 ordinal(0)
// CHECK-SAME: layout(#hal.executable.layout<push_constants = 0, sets = [
-// CHECK-SAME: #hal.descriptor_set.layout<0, bindings = [
-// CHECK-SAME: #hal.descriptor_set.binding<0, storage_buffer>,
-// CHECK-SAME: #hal.descriptor_set.binding<1, storage_buffer>,
-// CHECK-SAME: #hal.descriptor_set.binding<2, storage_buffer>
+// CHECK-SAME: <0, bindings = [
+// CHECK-SAME: <0, storage_buffer>,
+// CHECK-SAME: <1, storage_buffer>,
+// CHECK-SAME: <2, storage_buffer>
// CHECK: module attributes {vm.toplevel} {
// CHECK-NEXT: vm.module public @module {
// CHECK-NEXT: vm.func private @add_dispatch_0(
diff --git a/compiler/src/iree/compiler/InputConversion/MHLO/BroadcastingToLinalgPatterns.cpp b/compiler/src/iree/compiler/InputConversion/MHLO/BroadcastingToLinalgPatterns.cpp
index d47d085..0dbedb5 100644
--- a/compiler/src/iree/compiler/InputConversion/MHLO/BroadcastingToLinalgPatterns.cpp
+++ b/compiler/src/iree/compiler/InputConversion/MHLO/BroadcastingToLinalgPatterns.cpp
@@ -379,7 +379,7 @@
Operation *op, ArrayRef<Value> operands) override {
auto broadcastDimensions = llvm::cast<FromOpTy>(op).broadcast_dimensions();
if (broadcastDimensions &&
- !hlo::IsLegalNumpyRankedBroadcast(operands[0], operands[1],
+ !hlo::isLegalNumpyRankedBroadcast(operands[0], operands[1],
*broadcastDimensions)) {
return failure();
}
@@ -412,7 +412,7 @@
auto broadcastDimensions =
llvm::cast<chlo::BroadcastCompareOp>(op).broadcast_dimensions();
if (broadcastDimensions &&
- !hlo::IsLegalNumpyRankedBroadcast(operands[0], operands[1],
+ !hlo::isLegalNumpyRankedBroadcast(operands[0], operands[1],
*broadcastDimensions)) {
return failure();
}
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/ForeachThreadToSequentialFor.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/ForeachThreadToSequentialFor.cpp
index b880a48..e62aa9d 100644
--- a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/ForeachThreadToSequentialFor.cpp
+++ b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/ForeachThreadToSequentialFor.cpp
@@ -25,9 +25,10 @@
namespace {
SmallVector<Value> getValuesToYield(scf::PerformConcurrentlyOp op) {
- return llvm::to_vector(llvm::map_range(op.yieldingOps(), [](Operation &op) {
- return cast<scf::ParallelInsertSliceOp>(&op).getDest();
- }));
+ return llvm::to_vector(
+ llvm::map_range(op.getYieldingOps(), [](Operation &op) {
+ return cast<tensor::ParallelInsertSliceOp>(&op).getDest();
+ }));
}
} // namespace
@@ -79,9 +80,9 @@
// Create sequential insertSlice ops.
SmallVector<Value> toYield;
rewriter.setInsertionPoint(performConcurrentlyOp);
- for (Operation &operation : performConcurrentlyOp.yieldingOps()) {
- scf::ParallelInsertSliceOp op =
- cast<scf::ParallelInsertSliceOp>(&operation);
+ for (Operation &operation : performConcurrentlyOp.getYieldingOps()) {
+ tensor::ParallelInsertSliceOp op =
+ cast<tensor::ParallelInsertSliceOp>(&operation);
toYield.push_back(rewriter.createOrFold<tensor::InsertSliceOp>(
loc, op.getSource(), bvm.lookup(op.getDest()), op.getMixedOffsets(),
op.getMixedSizes(), op.getMixedStrides()));
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/Utils.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/Utils.cpp
index 5d6cc9f..d76d4fa 100644
--- a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/Utils.cpp
+++ b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/Transforms/Utils.cpp
@@ -90,7 +90,7 @@
void mlir::iree_compiler::IREE::LinalgExt::createMatchingParallelSubsetInsertOp(
OpBuilder &b, Location loc, tensor::ExtractSliceOp subsetExtractOp,
Value source, Value dest) {
- b.create<scf::ParallelInsertSliceOp>(
+ b.create<tensor::ParallelInsertSliceOp>(
loc, source, dest, subsetExtractOp.getMixedOffsets(),
subsetExtractOp.getMixedSizes(), subsetExtractOp.getMixedStrides());
}
diff --git a/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/fuse-in-containing-op.mlir b/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/fuse-in-containing-op.mlir
index 1ae34bb..39bd7d5 100644
--- a/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/fuse-in-containing-op.mlir
+++ b/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/fuse-in-containing-op.mlir
@@ -28,7 +28,7 @@
// CHECK: %[[T2:.*]] = linalg.elemwise_unary ins(%[[T1]]
%7 = linalg.elemwise_unary ins(%6 : tensor<?xf32>) outs(%5 : tensor<?xf32>) -> tensor<?xf32>
scf.foreach_thread.perform_concurrently {
- scf.foreach_thread.parallel_insert_slice %7 into %arg2[%3] [%4] [1] : tensor<?xf32> into tensor<64xf32>
+ tensor.parallel_insert_slice %7 into %arg2[%3] [%4] [1] : tensor<?xf32> into tensor<64xf32>
}
}
func.return %2 : tensor<64xf32>
@@ -90,7 +90,7 @@
// CHECK: %[[T2:.*]] = linalg.elemwise_unary ins(%[[T1]]
%7 = linalg.elemwise_unary ins(%6 : tensor<?xf32>) outs(%5 : tensor<?xf32>) -> tensor<?xf32>
scf.foreach_thread.perform_concurrently {
- scf.foreach_thread.parallel_insert_slice %7 into %arg2[%3] [%4] [1] : tensor<?xf32> into tensor<?xf32>
+ tensor.parallel_insert_slice %7 into %arg2[%3] [%4] [1] : tensor<?xf32> into tensor<?xf32>
}
}
func.return %2 : tensor<?xf32>
diff --git a/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/fuse-operands.mlir b/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/fuse-operands.mlir
index d4fa91a..629fba5 100644
--- a/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/fuse-operands.mlir
+++ b/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/fuse-operands.mlir
@@ -32,7 +32,7 @@
// CHECK: %[[T4:.*]] = linalg.elemwise_unary ins(%[[T1]] {{.*}} outs(%[[T3]]
%8 = linalg.elemwise_unary ins(%6 : tensor<?xf32>) outs(%7 : tensor<?xf32>) -> tensor<?xf32>
scf.foreach_thread.perform_concurrently {
- scf.foreach_thread.parallel_insert_slice %8 into %arg2[%4] [%5] [1] : tensor<?xf32> into tensor<64xf32>
+ tensor.parallel_insert_slice %8 into %arg2[%4] [%5] [1] : tensor<?xf32> into tensor<64xf32>
}
}
func.return %3 : tensor<64xf32>
@@ -93,7 +93,7 @@
// CHECK: %[[T2:.*]] = linalg.elemwise_unary ins(%[[T1]]
%7 = linalg.elemwise_unary ins(%6 : tensor<?xf32>) outs(%5 : tensor<?xf32>) -> tensor<?xf32>
scf.foreach_thread.perform_concurrently {
- scf.foreach_thread.parallel_insert_slice %7 into %arg2[%3] [%4] [1] : tensor<?xf32> into tensor<?xf32>
+ tensor.parallel_insert_slice %7 into %arg2[%3] [%4] [1] : tensor<?xf32> into tensor<?xf32>
}
}
func.return %2 : tensor<?xf32>
diff --git a/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/tiling-to-foreach-thread-op.mlir b/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/tiling-to-foreach-thread-op.mlir
index 8e25198..437e006 100644
--- a/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/tiling-to-foreach-thread-op.mlir
+++ b/llvm-external-projects/iree-dialects/test/Dialect/iree_linalg_ext/tiling-to-foreach-thread-op.mlir
@@ -23,7 +23,7 @@
// CHECK-SAME: ins(%[[tA]], %[[tB]] : tensor<?x?xf32>, tensor<?x?xf32>)
// CHECK-SAME: outs(%[[tC]] : tensor<?x?xf32>) -> tensor<?x?xf32>
// CHECK-NEXT: scf.foreach_thread.perform_concurrently {
- // CHECK-NEXT: scf.foreach_thread.parallel_insert_slice %[[RES]] into %[[C]]{{.*}} :
+ // CHECK-NEXT: tensor.parallel_insert_slice %[[RES]] into %[[C]]{{.*}} :
// CHECK-SAME: tensor<?x?xf32> into tensor<?x?xf32>
// CHECK-NEXT: }
// CHECK-NEXT: } {thread_dim_mapping = [1, 0]}
diff --git a/tests/e2e/models/unidirectional_lstm.mlir b/tests/e2e/models/unidirectional_lstm.mlir
index adea1cb..cedbf61 100644
--- a/tests/e2e/models/unidirectional_lstm.mlir
+++ b/tests/e2e/models/unidirectional_lstm.mlir
@@ -129,13 +129,13 @@
%98 = "mhlo.reshape"(%42) : (tensor<i64>) -> tensor<1xi64>
%99 = "mhlo.convert"(%98) : (tensor<1xi64>) -> tensor<1xi32>
%100 = "mhlo.reshape"(%99) : (tensor<1xi32>) -> tensor<i32>
- %101 = "mhlo.dynamic-update-slice"(%53, %97, %100) : (tensor<5xi64>, tensor<1xi64>, tensor<i32>) -> tensor<5xi64>
+ %101 = "mhlo.dynamic_update_slice"(%53, %97, %100) : (tensor<5xi64>, tensor<1xi64>, tensor<i32>) -> tensor<5xi64>
%102 = "mhlo.reshape"(%84) : (tensor<1x10xf32>) -> tensor<1x1x10xf32>
%103 = "mhlo.reshape"(%99) : (tensor<1xi32>) -> tensor<i32>
- %104 = "mhlo.dynamic-update-slice"(%54, %102, %103, %cst_6, %cst_6) : (tensor<5x1x10xf32>, tensor<1x1x10xf32>, tensor<i32>, tensor<i32>, tensor<i32>) -> tensor<5x1x10xf32>
+ %104 = "mhlo.dynamic_update_slice"(%54, %102, %103, %cst_6, %cst_6) : (tensor<5x1x10xf32>, tensor<1x1x10xf32>, tensor<i32>, tensor<i32>, tensor<i32>) -> tensor<5x1x10xf32>
%105 = "mhlo.reshape"(%96) : (tensor<1x10xf32>) -> tensor<1x1x10xf32>
%106 = "mhlo.reshape"(%99) : (tensor<1xi32>) -> tensor<i32>
- %107 = "mhlo.dynamic-update-slice"(%55, %105, %106, %cst_6, %cst_6) : (tensor<5x1x10xf32>, tensor<1x1x10xf32>, tensor<i32>, tensor<i32>, tensor<i32>) -> tensor<5x1x10xf32>
+ %107 = "mhlo.dynamic_update_slice"(%55, %105, %106, %cst_6, %cst_6) : (tensor<5x1x10xf32>, tensor<1x1x10xf32>, tensor<i32>, tensor<i32>, tensor<i32>) -> tensor<5x1x10xf32>
cf.br ^bb1(%56, %43, %44, %45, %46, %47, %84, %96, %50, %51, %52, %101, %104, %107 : tensor<i64>, tensor<i64>, tensor<40xf32>, tensor<i64>, tensor<74x40xf32>, tensor<i64>, tensor<1x10xf32>, tensor<1x10xf32>, tensor<5x1x64xf32>, tensor<5x1x1xf32>, tensor<5x1x1xf32>, tensor<5xi64>, tensor<5x1x10xf32>, tensor<5x1x10xf32>)
^bb3(%108: tensor<i64>, %109: tensor<i64>, %110: tensor<1x10xf32>, %111: tensor<1x10xf32>, %112: tensor<5xi64>, %113: tensor<5x1x10xf32>, %114: tensor<5x1x10xf32>): // pred: ^bb1
return %108, %112, %113, %114, %109, %110, %111 : tensor<i64>, tensor<5xi64>, tensor<5x1x10xf32>, tensor<5x1x10xf32>, tensor<i64>, tensor<1x10xf32>, tensor<1x10xf32>
diff --git a/tests/e2e/xla_ops/dynamic_update_slice.mlir b/tests/e2e/xla_ops/dynamic_update_slice.mlir
index bfdfa05..7031373 100644
--- a/tests/e2e/xla_ops/dynamic_update_slice.mlir
+++ b/tests/e2e/xla_ops/dynamic_update_slice.mlir
@@ -2,7 +2,7 @@
%target = util.unfoldable_constant dense<2> : tensor<3x3xi32>
%update = util.unfoldable_constant dense<1> : tensor<2x2xi32>
%c0 = util.unfoldable_constant dense<0> : tensor<i32>
- %result = "mhlo.dynamic-update-slice"(%target, %update, %c0, %c0)
+ %result = "mhlo.dynamic_update_slice"(%target, %update, %c0, %c0)
: (tensor<3x3xi32>, tensor<2x2xi32>, tensor<i32>, tensor<i32>) -> tensor<3x3xi32>
check.expect_eq_const(%result, dense<[
[1, 1, 2],
@@ -16,7 +16,7 @@
%update = util.unfoldable_constant dense<1> : tensor<1x3xi32>
%c0 = util.unfoldable_constant dense<0> : tensor<i32>
%c1 = util.unfoldable_constant dense<1> : tensor<i32>
- %result = "mhlo.dynamic-update-slice"(%target, %update, %c1, %c0)
+ %result = "mhlo.dynamic_update_slice"(%target, %update, %c1, %c0)
: (tensor<3x3xi32>, tensor<1x3xi32>, tensor<i32>, tensor<i32>) -> tensor<3x3xi32>
check.expect_eq_const(%result, dense<[
[2, 2, 2],
@@ -29,7 +29,7 @@
%update = util.unfoldable_constant dense<2> : tensor<1xi32>
%target = mhlo.constant dense<1> : tensor<4xi32>
%index = mhlo.constant dense<0> : tensor<i32>
- %result = "mhlo.dynamic-update-slice"(%target, %update, %index) : (tensor<4xi32>, tensor<1xi32>, tensor<i32>) -> tensor<4xi32>
+ %result = "mhlo.dynamic_update_slice"(%target, %update, %index) : (tensor<4xi32>, tensor<1xi32>, tensor<i32>) -> tensor<4xi32>
check.expect_eq_const(%result, dense<[2, 1, 1, 1]> : tensor<4xi32>) : tensor<4xi32>
return
}
diff --git a/third_party/llvm-project b/third_party/llvm-project
index 52e559b..52203a5 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit 52e559b2c0e8d362ac7f81a31b65cb3a3aa66fda
+Subproject commit 52203a575295fcaa51ccfab15b00ce15f6f7408b
diff --git a/third_party/mlir-hlo b/third_party/mlir-hlo
index df41f42..7365a77 160000
--- a/third_party/mlir-hlo
+++ b/third_party/mlir-hlo
@@ -1 +1 @@
-Subproject commit df41f42da81a59c23391f994d4bced6afc95a505
+Subproject commit 7365a774d0c772e5a032ac0d36a1d532f17860f3
