Integrate LLVM at llvm/llvm-project@95995673d1ba
Updates LLVM usage to match
[95995673d1ba](https://github.com/llvm/llvm-project/commit/95995673d1ba)
PiperOrigin-RevId: 385228604
diff --git a/SUBMODULE_VERSIONS.txt b/SUBMODULE_VERSIONS.txt
index 4ba50ca..311a796 100644
--- a/SUBMODULE_VERSIONS.txt
+++ b/SUBMODULE_VERSIONS.txt
@@ -4,7 +4,7 @@
aa533abfd4232b01f9e57041d70114d5a77e6de0 third_party/googletest
88b845dee001723c4a0db1fe5477de735b6d3bb0 third_party/liburing
acd6f6f014c25e46363e718381e0b35205df2d83 third_party/libyaml
-d046fb62b7e7cd273b104fee0162725003411c00 third_party/llvm-project
+95995673d1babee5731146f45ad3ce79c32f06d4 third_party/llvm-project
5a6c67a5c32199e240cd8e5f8ad38b0821371031 third_party/mlir-hlo
4c7697dbe973ed01ae6fbec37d186ebd05982e1f third_party/pybind11
2e1b5fb39ebc2ef4cb77005f8267e4f3a6241ba1 third_party/spirv_cross
diff --git a/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp b/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
index 403e436..f6c3394 100644
--- a/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
+++ b/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
@@ -855,7 +855,7 @@
resultBuffer.getType().cast<MemRefType>().getMemorySpaceAsInt());
using ReverseReshapeOpTy = typename std::conditional<
std::is_same<TensorReshapeOpTy, linalg::TensorCollapseShapeOp>::value,
- linalg::ExpandShapeOp, linalg::CollapseShapeOp>::type;
+ memref::ExpandShapeOp, memref::CollapseShapeOp>::type;
return b.create<ReverseReshapeOpTy>(reshapeOp.getLoc(), memrefType,
resultBuffer, reshapeOp.reassociation());
}
@@ -1004,7 +1004,7 @@
resultTensorType, {}, inputBufferType.getMemorySpaceAsInt());
using ReshapeOpTy = typename std::conditional<
std::is_same<TensorReshapeOpTy, linalg::TensorCollapseShapeOp>::value,
- linalg::CollapseShapeOp, linalg::ExpandShapeOp>::type;
+ memref::CollapseShapeOp, memref::ExpandShapeOp>::type;
Value bufferReshape = b.create<ReshapeOpTy>(loc, reshapeResultType,
inputBuffer, op.reassociation());
return bufferReshape;
diff --git a/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
index f954db2..f70ea49 100644
--- a/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
+++ b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
@@ -655,7 +655,7 @@
// CHECK: func @reshape_simple()
// CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0
// CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0
-// CHECK: %[[RESHAPE:.+]] = linalg.expand_shape %[[ARG0]] {{\[}}[0, 1]]
+// CHECK: %[[RESHAPE:.+]] = memref.expand_shape %[[ARG0]] {{\[}}[0, 1]]
// CHECK: linalg.copy(%[[RESHAPE]], %[[RET0]])
// -----
@@ -690,7 +690,7 @@
// CHECK: %[[C0:.+]] = constant 0
// CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0[%[[C0]]] : memref<12xi32>
// CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0[%[[C0]]] : memref<3x4xi32>
-// CHECK: %[[RESHAPE:.+]] = linalg.expand_shape %[[ARG0]] {{\[}}[0, 1]]
+// CHECK: %[[RESHAPE:.+]] = memref.expand_shape %[[ARG0]] {{\[}}[0, 1]]
// CHECK: linalg.generic
// CHECK-SAME: ins(%[[RESHAPE]] : memref<3x4xi32>)
// CHECK-SAME: outs(%[[RET0]] : memref<3x4xi32>)
@@ -731,7 +731,7 @@
// CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0[%[[C0]]] : memref<12xi32>
// CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0[%[[C0]]] : memref<3x4xi32>
// CHECK-DAG: %[[RET1:.+]] = hal.interface.binding.subspan @io::@ret1[%[[C0]]] : memref<3x4xi32>
-// CHECK: %[[RESHAPE:.+]] = linalg.expand_shape %[[ARG0]] {{\[}}[0, 1]]
+// CHECK: %[[RESHAPE:.+]] = memref.expand_shape %[[ARG0]] {{\[}}[0, 1]]
// CHECK: linalg.generic
// CHECK-SAME: ins(%[[RESHAPE]] : memref<3x4xi32>)
// CHECK-SAME: outs(%[[RET0]] : memref<3x4xi32>)
@@ -769,7 +769,7 @@
// CHECK: %[[C0:.+]] = constant 0
// CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0[%[[C0]]] : memref<3x4xi32>
// CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0[%[[C0]]] : memref<12xi32>
-// CHECK: %[[RESHAPE:.+]] = linalg.expand_shape %[[RET0]] {{\[}}[0, 1]]
+// CHECK: %[[RESHAPE:.+]] = memref.expand_shape %[[RET0]] {{\[}}[0, 1]]
// CHECK: linalg.generic
// CHECK-SAME: ins(%[[ARG0]] : memref<3x4xi32>)
// CHECK-SAME: outs(%[[RESHAPE]] : memref<3x4xi32>)
@@ -819,7 +819,7 @@
// CHECK-LABEL: func @dot_general_lowering()
// CHECK-DAG: %[[LHS:.+]] = hal.interface.binding.subspan @io::@arg0
// CHECK-DAG: %[[RHS:.+]] = hal.interface.binding.subspan @io::@arg1
-// CHECK-DAG: %[[RESHAPE_LHS:.+]] = linalg.collapse_shape %[[LHS]]
+// CHECK-DAG: %[[RESHAPE_LHS:.+]] = memref.collapse_shape %[[LHS]]
// CHECK-DAG: %[[RETURN:.+]] = hal.interface.binding.subspan @io::@ret0
// CHECK: scf.for %[[IV0:.+]] = {{.+}} {
// CHECK: scf.for %[[IV1:.+]] = {{.+}} {
@@ -1097,7 +1097,7 @@
// CHECK-DAG: %[[RHS:.+]] = memref.buffer_cast %[[CONSTANT]]
// CHECK-DAG: %[[LHS_INPUT:.+]] = hal.interface.binding.subspan @io::@arg0[%{{.+}}] : memref<1x5x3x1xf32>
// CHECK-DAG: %[[RETURN:.+]] = hal.interface.binding.subspan @io::@ret0[%{{.+}}] : memref<5x5xf32>
-// CHECK: %[[LHS:.+]] = linalg.collapse_shape %[[LHS_INPUT]]
+// CHECK: %[[LHS:.+]] = memref.collapse_shape %[[LHS_INPUT]]
// CHECK: scf.for %[[IV0:.+]] =
// CHECK: scf.for %[[IV1:.+]] =
// CHECK-DAG: %[[LHS_SUBVIEW:.+]] = memref.subview %[[LHS]][%[[IV0]], 0]
@@ -1279,7 +1279,7 @@
// CHECK-LABEL: func @reshape_read_only
// CHECK-DAG: %[[INPUT:.+]] = hal.interface.binding.subspan @io::@ro0
// CHECK-DAG: %[[OUTPUT:.+]] = hal.interface.binding.subspan @io::@wo0
-// CHECK: %[[RESHAPE:.+]] = linalg.collapse_shape %[[INPUT]]
+// CHECK: %[[RESHAPE:.+]] = memref.collapse_shape %[[INPUT]]
// CHECK: linalg.generic
// CHECK-SAME: ins(%[[RESHAPE]] : memref<?xf32>)
// CHECK-SAME: outs(%[[OUTPUT]] : memref<?xf32>)
@@ -1916,9 +1916,9 @@
// CHECK: linalg.generic
// CHECK-SAME: ins(%[[ARG0_SV]]
// CHECK-SAME: outs(%[[ALLOC_ARG0]]
-// CHECK-DAG: %[[ALLOC_ARG0_RESHAPE:.+]] = linalg.collapse_shape %[[ALLOC_ARG0]]
-// CHECK-DAG: %[[ALLOC_ARG1_RESHAPE:.+]] = linalg.collapse_shape %[[ALLOC_ARG1]]
-// CHECK-DAG: %[[ALLOC_RET0_RESHAPE:.+]] = linalg.collapse_shape %[[ALLOC_RET0]]
+// CHECK-DAG: %[[ALLOC_ARG0_RESHAPE:.+]] = memref.collapse_shape %[[ALLOC_ARG0]]
+// CHECK-DAG: %[[ALLOC_ARG1_RESHAPE:.+]] = memref.collapse_shape %[[ALLOC_ARG1]]
+// CHECK-DAG: %[[ALLOC_RET0_RESHAPE:.+]] = memref.collapse_shape %[[ALLOC_RET0]]
// CHECK: linalg.matmul
// CHECK-SAME: ins(%[[ALLOC_ARG0_RESHAPE]], %[[ALLOC_ARG1_RESHAPE]]
// CHECK-SAME: outs(%[[ALLOC_RET0_RESHAPE]]
diff --git a/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp b/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
index 6fc6a65..f8f1753 100644
--- a/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
@@ -281,9 +281,11 @@
}
Value getIndexValue(Location loc, int64_t value, OpBuilder &builder) {
- return builder.createOrFold<LLVM::DialectCastOp>(
- loc, typeConverter->convertType(builder.getIndexType()),
+ SmallVector<Value> folded;
+ builder.createOrFold<UnrealizedConversionCastOp>(
+ folded, loc, typeConverter->convertType(builder.getIndexType()),
builder.createOrFold<ConstantIndexOp>(loc, value));
+ return folded.front();
}
Value castValueToType(Location loc, Value value, Type resultType,
@@ -731,6 +733,7 @@
target.addIllegalDialect<ShapeDialect, StandardOpsDialect, IREEDialect,
IREE::HAL::HALDialect, math::MathDialect,
tosa::TosaDialect>();
+ target.addIllegalOp<UnrealizedConversionCastOp>();
// Don't apply patterns to private function (e.g num_workgroups func).
target.addDynamicallyLegalOp<FuncOp>([&](FuncOp funcOp) {
diff --git a/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp b/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp
index b1affc9..002299f 100644
--- a/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp
+++ b/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp
@@ -315,7 +315,7 @@
/// !spv.array.
/// - unrealized_conversion_cast with the same source and target type.
patterns.insert<
- FoldAsNoOp<linalg::CollapseShapeOp>, FoldAsNoOp<linalg::ExpandShapeOp>,
+ FoldAsNoOp<memref::CollapseShapeOp>, FoldAsNoOp<memref::ExpandShapeOp>,
FoldAsNoOp<memref::BufferCastOp>, RemoveIdentityConversionCast>(
typeConverter, context);
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVConvertToGPU.cpp b/iree/compiler/Codegen/SPIRV/SPIRVConvertToGPU.cpp
index 8899705..e154864 100644
--- a/iree/compiler/Codegen/SPIRV/SPIRVConvertToGPU.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVConvertToGPU.cpp
@@ -653,7 +653,7 @@
// Reshape ops are treated legal since they just change the way the underlying
// buffer is viewed. These are legalized downstream. They become no ops when
// lowering to SPIR-V since the SPIR-V code uses linearized arrays.
- target.addLegalOp<linalg::CollapseShapeOp, linalg::ExpandShapeOp>();
+ target.addLegalOp<memref::CollapseShapeOp, memref::ExpandShapeOp>();
// Let the rest fall through.
target.markUnknownOpDynamicallyLegal([](Operation *) { return true; });
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
index 6a3cf77..4c3d839 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
@@ -109,8 +109,8 @@
%0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<1x113x113x96xf32>
%1 = hal.interface.binding.subspan @io::@arg1[%c0] : memref<3x3x1x96xf32>
%2 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<1x56x56x96xf32>
- %3 = linalg.collapse_shape %1 [[0, 1, 2, 3]] : memref<3x3x1x96xf32> into memref<864xf32>
- %4 = linalg.expand_shape %3 [[0, 1, 2]] : memref<864xf32> into memref<3x3x96xf32>
+ %3 = memref.collapse_shape %1 [[0, 1, 2, 3]] : memref<3x3x1x96xf32> into memref<864xf32>
+ %4 = memref.expand_shape %3 [[0, 1, 2]] : memref<864xf32> into memref<3x3x96xf32>
%workgroup_size_x = hal.interface.workgroup.size[0] : index
%workgroup_size_y = hal.interface.workgroup.size[1] : index
%workgroup_size_z = hal.interface.workgroup.size[2] : index
diff --git a/iree/compiler/Dialect/Modules/VMVX/Conversion/StandardToVMVX/ConvertStandardToVMVX.cpp b/iree/compiler/Dialect/Modules/VMVX/Conversion/StandardToVMVX/ConvertStandardToVMVX.cpp
index 96e8fe1..40b8914 100644
--- a/iree/compiler/Dialect/Modules/VMVX/Conversion/StandardToVMVX/ConvertStandardToVMVX.cpp
+++ b/iree/compiler/Dialect/Modules/VMVX/Conversion/StandardToVMVX/ConvertStandardToVMVX.cpp
@@ -68,8 +68,8 @@
OwningRewritePatternList &patterns,
TypeConverter &typeConverter) {
// We type/shape erase memrefs as we lower so there is no need for reshapes.
- patterns.insert<FoldAsNoOp<linalg::CollapseShapeOp>>(typeConverter, context);
- patterns.insert<FoldAsNoOp<linalg::ExpandShapeOp>>(typeConverter, context);
+ patterns.insert<FoldAsNoOp<memref::CollapseShapeOp>>(typeConverter, context);
+ patterns.insert<FoldAsNoOp<memref::ExpandShapeOp>>(typeConverter, context);
patterns.insert<RemoveIdentityConversionCast>(typeConverter, context);
}
diff --git a/iree/tools/BUILD b/iree/tools/BUILD
index 612ef75..4ca84ad 100644
--- a/iree/tools/BUILD
+++ b/iree/tools/BUILD
@@ -145,6 +145,7 @@
"@llvm-project//mlir:LinalgToSPIRV",
"@llvm-project//mlir:LinalgTransforms",
"@llvm-project//mlir:MathDialect",
+ "@llvm-project//mlir:MemRefDialect",
"@llvm-project//mlir:QuantOps",
"@llvm-project//mlir:SCFDialect",
"@llvm-project//mlir:SCFToGPUPass",
diff --git a/third_party/llvm-project b/third_party/llvm-project
index d046fb6..9599567 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit d046fb62b7e7cd273b104fee0162725003411c00
+Subproject commit 95995673d1babee5731146f45ad3ce79c32f06d4
