Merge branch 'google' into main-to-google
diff --git a/SUBMODULE_VERSIONS.txt b/SUBMODULE_VERSIONS.txt
index 62fb1ca..5d408df 100644
--- a/SUBMODULE_VERSIONS.txt
+++ b/SUBMODULE_VERSIONS.txt
@@ -4,13 +4,13 @@
aa533abfd4232b01f9e57041d70114d5a77e6de0 third_party/googletest
88b845dee001723c4a0db1fe5477de735b6d3bb0 third_party/liburing
acd6f6f014c25e46363e718381e0b35205df2d83 third_party/libyaml
-d046fb62b7e7cd273b104fee0162725003411c00 third_party/llvm-project
-5a6c67a5c32199e240cd8e5f8ad38b0821371031 third_party/mlir-hlo
+5994201c8e4c69518a10e143d09973b77886c90f third_party/llvm-project
+1b6e71edd1df812640d8453814afccb1bc12e170 third_party/mlir-hlo
4c7697dbe973ed01ae6fbec37d186ebd05982e1f third_party/pybind11
2e1b5fb39ebc2ef4cb77005f8267e4f3a6241ba1 third_party/spirv_cross
f5417a4b6633c3217c9a1bc2f0c70b1454975ba7 third_party/spirv_headers
b42009b3b9d4ca35bc703f5310eedc74f584be58 third_party/stblib
-0e956f76c739425e45a3e1d15e0d84d782011d86 third_party/tensorflow
+951452b6150ee86e2562aaf46a4bcceb4aa99697 third_party/tensorflow
50f7deb1a389bd3785c12fbe0be74128343f11f7 third_party/tracy
9d10a96f2d57c3c37e167f2e73c9a31ac2e51fa5 third_party/vulkan_headers
8d4a9e9174a9c6ad6a3a3ae981b915ef13fc12c4 third_party/vulkan_memory_allocator
diff --git a/bindings/python/iree/runtime/function.py b/bindings/python/iree/runtime/function.py
index 57537e7..36a5280 100644
--- a/bindings/python/iree/runtime/function.py
+++ b/bindings/python/iree/runtime/function.py
@@ -338,6 +338,7 @@
"i32": np.int32,
"i64": np.int64,
"f64": np.float64,
+ "i16": np.int16,
"i1": np.bool_,
}
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/LLVMCPU/LLVMCPUTilePadAndVectorize.cpp b/iree/compiler/Codegen/LLVMCPU/LLVMCPUTilePadAndVectorize.cpp
index cd8fdfe..2fd40fc 100644
--- a/iree/compiler/Codegen/LLVMCPU/LLVMCPUTilePadAndVectorize.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/LLVMCPUTilePadAndVectorize.cpp
@@ -170,7 +170,7 @@
.insert<ContractionOpToOuterProductOpLowering,
ContractionOpToMatmulOpLowering, ContractionOpLowering>(
vectorTransformsOptions, context);
- vector::populateVectorTransferLoweringPatterns(
+ vector::populateVectorTransferPermutationMapLoweringPatterns(
vectorContractLoweringPatterns);
if (failed(applyPatternsAndFoldGreedily(
funcOp, std::move(vectorContractLoweringPatterns)))) {
diff --git a/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorization.cpp b/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorization.cpp
index 7969658..aed1951 100644
--- a/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorization.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorization.cpp
@@ -263,7 +263,7 @@
.insert<ContractionOpToOuterProductOpLowering,
ContractionOpToMatmulOpLowering, ContractionOpLowering>(
vectorTransformsOptions, context);
- vector::populateVectorTransferLoweringPatterns(
+ vector::populateVectorTransferPermutationMapLoweringPatterns(
vectorContractLoweringPatterns);
if (failed(applyPatternsAndFoldGreedily(
funcOp, std::move(vectorContractLoweringPatterns)))) {
diff --git a/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp b/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp
index f04a19f..7544522 100644
--- a/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp
@@ -59,6 +59,7 @@
vector::VectorTransformsOptions().setVectorTransformsOptions(
vector::VectorContractLowering::OuterProduct));
mlir::vector::populateVectorTransposeLoweringPatterns(patterns);
+ mlir::vector::populateVectorTransferLoweringPatterns(patterns);
(void)applyPatternsAndFoldGreedily(m, std::move(patterns));
}
{
diff --git a/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp b/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
index 00b2990..b6cf257 100644
--- a/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
@@ -59,6 +59,7 @@
vector::VectorTransformsOptions().setVectorTransformsOptions(
vector::VectorContractLowering::OuterProduct));
mlir::vector::populateVectorTransposeLoweringPatterns(patterns);
+ vector::populateVectorTransferLoweringPatterns(patterns);
(void)applyPatternsAndFoldGreedily(m, std::move(patterns));
}
{
diff --git a/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp b/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
index b3276ea..1f62944 100644
--- a/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
@@ -160,10 +160,8 @@
linalg::DepthwiseConvInputNHWCFilterHWCOp,
linalg::ConvInputNHWCFilterHWCFOp,
linalg::DepthwiseConvInputNHWCFilterHWCFOp,
- linalg::DepthwiseConvInputNHWCFilterHWCOp,
- linalg::PoolingNHWCMaxI8Op, linalg::PoolingNHWCMaxI16Op,
- linalg::PoolingNHWCMaxI32Op, linalg::PoolingNHWCMaxFOp,
- linalg::PoolingNHWCMinFOp, linalg::PoolingNHWCSumFOp>(
+ linalg::DepthwiseConvInputNHWCFilterHWCOp, linalg::PoolingNhwcMaxOp,
+ linalg::PoolingNhwcMinOp, linalg::PoolingNhwcSumOp>(
op.getOperation())) {
rootOperation = op;
break;
diff --git a/iree/compiler/Codegen/LLVMGPU/LLVMGPUVectorization.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPUVectorization.cpp
index 0df5145..e1a035e 100644
--- a/iree/compiler/Codegen/LLVMGPU/LLVMGPUVectorization.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPUVectorization.cpp
@@ -103,7 +103,8 @@
RewritePatternSet lowerTransferOpPatterns(funcOp.getContext());
vector::populateVectorToVectorCanonicalizationPatterns(
lowerTransferOpPatterns);
- vector::populateVectorTransferLoweringPatterns(lowerTransferOpPatterns);
+ vector::populateVectorTransferPermutationMapLoweringPatterns(
+ lowerTransferOpPatterns);
(void)applyPatternsAndFoldGreedily(funcOp,
std::move(lowerTransferOpPatterns));
}
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/KernelDispatchUtils.cpp b/iree/compiler/Codegen/SPIRV/KernelDispatchUtils.cpp
index 1ab31ca..98b065c 100644
--- a/iree/compiler/Codegen/SPIRV/KernelDispatchUtils.cpp
+++ b/iree/compiler/Codegen/SPIRV/KernelDispatchUtils.cpp
@@ -750,12 +750,9 @@
DISPATCH(linalg::ConvInputNHWCFilterHWCFOp)
DISPATCH(linalg::ConvInputNDHWCFilterDHWCFOp)
DISPATCH(linalg::MatmulOp)
- DISPATCH(linalg::PoolingNHWCMaxI8Op)
- DISPATCH(linalg::PoolingNHWCMaxI16Op)
- DISPATCH(linalg::PoolingNHWCMaxI32Op)
- DISPATCH(linalg::PoolingNHWCMaxFOp)
- DISPATCH(linalg::PoolingNHWCMinFOp)
- DISPATCH(linalg::PoolingNHWCSumFOp)
+ DISPATCH(linalg::PoolingNhwcMaxOp)
+ DISPATCH(linalg::PoolingNhwcMinOp)
+ DISPATCH(linalg::PoolingNhwcSumOp)
}
// Any generic operations found are made the root if no other op is the root
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/SPIRVTileAndVectorize.cpp b/iree/compiler/Codegen/SPIRV/SPIRVTileAndVectorize.cpp
index 8fefc5d..04cc42b 100644
--- a/iree/compiler/Codegen/SPIRV/SPIRVTileAndVectorize.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVTileAndVectorize.cpp
@@ -285,9 +285,9 @@
linalg::LinalgTilingPattern<linalg::ConvInputNDHWCFilterDHWCFOp>,
linalg::LinalgTilingPattern<linalg::DepthwiseConvInputNHWCFilterHWCFOp>,
linalg::LinalgTilingPattern<linalg::GenericOp>,
- linalg::LinalgTilingPattern<linalg::PoolingNHWCMaxFOp>,
- linalg::LinalgTilingPattern<linalg::PoolingNHWCMinFOp>,
- linalg::LinalgTilingPattern<linalg::PoolingNHWCSumFOp>>(
+ linalg::LinalgTilingPattern<linalg::PoolingNhwcMaxOp>,
+ linalg::LinalgTilingPattern<linalg::PoolingNhwcMinOp>,
+ linalg::LinalgTilingPattern<linalg::PoolingNhwcSumOp>>(
context, tilingOptions,
getLinalgMatchAndReplaceMarker(
{getWorkgroupMemoryMarker(), getWorkgroupMarker()},
@@ -421,8 +421,18 @@
std::move(vectorUnrollPatterns));
}
{
+ linalg::hoistRedundantVectorTransfers(funcOp);
+
+ LLVM_DEBUG({
+ llvm::dbgs() << "--- After hoisting vector transfers ---\n";
+ funcOp.print(llvm::dbgs(), OpPrintingFlags().useLocalScope());
+ llvm::dbgs() << "\n\n";
+ });
+ }
+ {
RewritePatternSet canonicalizationPatterns2(funcOp.getContext());
- vector::populateVectorTransferLoweringPatterns(canonicalizationPatterns2);
+ vector::populateVectorTransferPermutationMapLoweringPatterns(
+ canonicalizationPatterns2);
(void)applyPatternsAndFoldGreedily(funcOp,
std::move(canonicalizationPatterns2));
@@ -453,16 +463,6 @@
llvm::dbgs() << "\n\n";
});
}
-
- {
- linalg::hoistRedundantVectorTransfers(funcOp);
-
- LLVM_DEBUG({
- llvm::dbgs() << "--- After hoisting vector transfers ---\n";
- funcOp.print(llvm::dbgs(), OpPrintingFlags().useLocalScope());
- llvm::dbgs() << "\n\n";
- });
- }
}
//====---------------------------------------------------------------------===//
@@ -722,9 +722,9 @@
LowerToLoops<linalg::DepthwiseConvInputNHWCFilterHWCOp>,
LowerToLoops<linalg::FillOp>, LowerToLoops<linalg::GenericOp>,
LowerToLoops<linalg::MatmulOp>,
- LowerToLoops<linalg::PoolingNHWCMaxFOp>,
- LowerToLoops<linalg::PoolingNHWCMinFOp>,
- LowerToLoops<linalg::PoolingNHWCSumFOp>>(context);
+ LowerToLoops<linalg::PoolingNhwcMaxOp>,
+ LowerToLoops<linalg::PoolingNhwcMinOp>,
+ LowerToLoops<linalg::PoolingNhwcSumOp>>(context);
(void)applyPatternsAndFoldGreedily(funcOp, std::move(patterns));
}
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/Flow/Transforms/DispatchLinalgOnTensors.cpp b/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
index 92fa3e1..759b5a1 100644
--- a/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
@@ -224,9 +224,8 @@
return isa<linalg::ConvInputNHWCFilterHWCFOp,
linalg::DepthwiseConvInputNHWCFilterHWCOp,
linalg::DepthwiseConvInputNHWCFilterHWCFOp,
- linalg::PoolingNHWCMaxI8Op, linalg::PoolingNHWCMaxI16Op,
- linalg::PoolingNHWCMaxI32Op, linalg::PoolingNHWCSumFOp,
- linalg::PoolingNHWCMaxFOp, linalg::PoolingNHWCMinFOp>(op);
+ linalg::PoolingNhwcSumOp, linalg::PoolingNhwcMaxOp,
+ linalg::PoolingNhwcMinOp>(op);
}
static bool isAlwaysClonedIntoDispatchOp(Operation *op) {
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..5994201 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit d046fb62b7e7cd273b104fee0162725003411c00
+Subproject commit 5994201c8e4c69518a10e143d09973b77886c90f
diff --git a/third_party/mlir-hlo b/third_party/mlir-hlo
index 5a6c67a..1b6e71e 160000
--- a/third_party/mlir-hlo
+++ b/third_party/mlir-hlo
@@ -1 +1 @@
-Subproject commit 5a6c67a5c32199e240cd8e5f8ad38b0821371031
+Subproject commit 1b6e71edd1df812640d8453814afccb1bc12e170
diff --git a/third_party/tensorflow b/third_party/tensorflow
index 0e956f7..951452b 160000
--- a/third_party/tensorflow
+++ b/third_party/tensorflow
@@ -1 +1 @@
-Subproject commit 0e956f76c739425e45a3e1d15e0d84d782011d86
+Subproject commit 951452b6150ee86e2562aaf46a4bcceb4aa99697
