diff --git a/SUBMODULE_VERSIONS.txt b/SUBMODULE_VERSIONS.txt
index 08ffcc2..7b22a43 100644
--- a/SUBMODULE_VERSIONS.txt
+++ b/SUBMODULE_VERSIONS.txt
@@ -4,7 +4,7 @@
 aa533abfd4232b01f9e57041d70114d5a77e6de0 third_party/googletest
 88b845dee001723c4a0db1fe5477de735b6d3bb0 third_party/liburing
 f8f760f7387d2cc56a2fc7b1be313a3bf3f7f58c third_party/libyaml
-def8b952ebc00b0ad0fa4196ae27b11a385087ad third_party/llvm-project
+c6a8bec4c578a7c12e4458b161fce7b1704804a2 third_party/llvm-project
 8c636d9692e2a50eb03d1e0a9809ffde90dbd2c2 third_party/mlir-hlo
 3f701faace7addc75d16dea8a6cd769fa5b3f260 third_party/musl
 59aa99860c60bd171b9565e9920f125fdb749267 third_party/pybind11
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/DirectLoweringPatterns.cpp b/integrations/tensorflow/iree_tf_compiler/TF/DirectLoweringPatterns.cpp
index 87fd143..7cdeb83 100644
--- a/integrations/tensorflow/iree_tf_compiler/TF/DirectLoweringPatterns.cpp
+++ b/integrations/tensorflow/iree_tf_compiler/TF/DirectLoweringPatterns.cpp
@@ -47,7 +47,7 @@
           op, "could not compute reassociation indices");
     }
 
-    rewriter.replaceOpWithNewOp<linalg::TensorCollapseShapeOp>(
+    rewriter.replaceOpWithNewOp<tensor::CollapseShapeOp>(
         op, resultType, op.input(), *reassociationIndices);
     return success();
   }
@@ -99,7 +99,7 @@
           op, "could not compute reassociation indices");
     }
 
-    rewriter.replaceOpWithNewOp<linalg::TensorExpandShapeOp>(
+    rewriter.replaceOpWithNewOp<tensor::ExpandShapeOp>(
         op, expandedType, op.input(), *reassociationIndices);
     return success();
   }
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/test/direct_lowering.mlir b/integrations/tensorflow/iree_tf_compiler/TF/test/direct_lowering.mlir
index 5316387..40c241c 100644
--- a/integrations/tensorflow/iree_tf_compiler/TF/test/direct_lowering.mlir
+++ b/integrations/tensorflow/iree_tf_compiler/TF/test/direct_lowering.mlir
@@ -2,7 +2,7 @@
 
 // CHECK-LABEL: @expand_dims
 func @expand_dims(%arg0: tensor<?x?x?xf32>) -> tensor<?x?x?x1xf32> {
-  // CHECK: %[[R:.*]] = linalg.tensor_expand_shape %arg0 {{\[}}[0], [1], [2, 3]] : tensor<?x?x?xf32> into tensor<?x?x?x1xf32>
+  // CHECK: %[[R:.*]] = tensor.expand_shape %arg0 {{\[}}[0], [1], [2, 3]] : tensor<?x?x?xf32> into tensor<?x?x?x1xf32>
   // CHECK: return %[[R]]
   %axis = "tf.Const"() {value = dense<3> : tensor<i32>} : () -> (tensor<i32>)
   %0 = "tf.ExpandDims"(%arg0, %axis) : (tensor<?x?x?xf32>, tensor<i32>) -> (tensor<?x?x?x1xf32>)
@@ -23,7 +23,7 @@
 // -----
 // CHECK-LABEL: @squeeze
 func @squeeze(%arg0 : tensor<?x1x1x1001xf32>) -> tensor<?x1001xf32> {
-  // CHECK: %[[R:.*]] = linalg.tensor_collapse_shape %arg0 {{\[}}[0], [1, 2, 3]] : tensor<?x1x1x1001xf32> into tensor<?x1001xf32>
+  // CHECK: %[[R:.*]] = tensor.collapse_shape %arg0 {{\[}}[0], [1, 2, 3]] : tensor<?x1x1x1001xf32> into tensor<?x1001xf32>
   // CHECK: return %[[R]]
   %0 = "tf.Squeeze"(%arg0) {device = "", squeeze_dims = [1, 2]} : (tensor<?x1x1x1001xf32>) -> tensor<?x1001xf32>
   return %0 : tensor<?x1001xf32>
diff --git a/iree/compiler/Codegen/Common/BufferizationAnalysis.cpp b/iree/compiler/Codegen/Common/BufferizationAnalysis.cpp
index 8cdce87..7664a5f 100644
--- a/iree/compiler/Codegen/Common/BufferizationAnalysis.cpp
+++ b/iree/compiler/Codegen/Common/BufferizationAnalysis.cpp
@@ -56,8 +56,8 @@
   // TODO(ravishankarm): Maybe this is too aggressive, might have to switch this
   // to have a white-list instead of blacklist.
   for (Operation *user : op->getUsers()) {
-    if (isa<IREE::Flow::DispatchTensorStoreOp, linalg::TensorCollapseShapeOp,
-            linalg::TensorExpandShapeOp>(user)) {
+    if (isa<IREE::Flow::DispatchTensorStoreOp, tensor::CollapseShapeOp,
+            tensor::ExpandShapeOp>(user)) {
       return false;
     }
   }
@@ -82,7 +82,7 @@
   return TypeSwitch<Operation *, bool>(definingOp)
       .Case<arith::ConstantOp>(
           [&](arith::ConstantOp constantOp) { return true; })
-      .Case<linalg::TensorCollapseShapeOp, linalg::TensorExpandShapeOp>(
+      .Case<tensor::CollapseShapeOp, tensor::ExpandShapeOp>(
           [&](auto op) { return isFromReadOnlyTensor(op.src(), plan); })
       .Case<tensor::ExtractSliceOp>([&](tensor::ExtractSliceOp sliceOp) {
         return isFromReadOnlyTensor(sliceOp.source(), plan);
@@ -159,7 +159,7 @@
       // reshapes not working well together, but there is no comment about why
       // this was added with the change that added this.
       Operation *op = v.getDefiningOp();
-      if (op && isa<linalg::TensorCollapseShapeOp, linalg::TensorExpandShapeOp>(
+      if (op && isa<tensor::CollapseShapeOp, tensor::ExpandShapeOp>(
                     v.getDefiningOp())) {
         return false;
       }
@@ -519,7 +519,7 @@
             [&](IREE::LinalgExt::LinalgExtOp linalgExtOp) {
               return analyseLinalgExtOps(linalgExtOp, plan);
             })
-        .Case<linalg::TensorCollapseShapeOp, linalg::TensorExpandShapeOp>(
+        .Case<tensor::CollapseShapeOp, tensor::ExpandShapeOp>(
             [&](auto reshapeOp) {
               return analyseSingleOperandResultOp(reshapeOp.src(),
                                                   reshapeOp.result(), plan);
diff --git a/iree/compiler/Codegen/Common/CleanupBufferAllocViewPass.cpp b/iree/compiler/Codegen/Common/CleanupBufferAllocViewPass.cpp
index 5484102..2d8b899 100644
--- a/iree/compiler/Codegen/Common/CleanupBufferAllocViewPass.cpp
+++ b/iree/compiler/Codegen/Common/CleanupBufferAllocViewPass.cpp
@@ -117,10 +117,9 @@
     : public CleanupBufferAllocViewBase<CleanupBufferAllocViewPass> {
   void runOnOperation() override {
     OwningRewritePatternList patterns(&getContext());
-    patterns.insert<
-        FoldReshapeIntoInterfaceTensorLoad<linalg::TensorCollapseShapeOp>,
-        FoldReshapeIntoInterfaceTensorLoad<linalg::TensorExpandShapeOp>,
-        RemoveDeadMemAllocs>(&getContext());
+    patterns.insert<FoldReshapeIntoInterfaceTensorLoad<tensor::CollapseShapeOp>,
+                    FoldReshapeIntoInterfaceTensorLoad<tensor::ExpandShapeOp>,
+                    RemoveDeadMemAllocs>(&getContext());
     if (failed(applyPatternsAndFoldGreedily(getOperation(),
                                             std::move(patterns)))) {
       return signalPassFailure();
diff --git a/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp b/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp
index 1ca48ce..e2d2f40 100644
--- a/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp
+++ b/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp
@@ -118,8 +118,8 @@
 static Value getReverseOfReshapeOp(OpBuilder &b, TensorReshapeOpTy reshapeOp,
                                    Value resultBuffer) {
   using ReverseReshapeOpTy = typename std::conditional<
-      std::is_same<TensorReshapeOpTy, linalg::TensorCollapseShapeOp>::value,
-      linalg::TensorExpandShapeOp, linalg::TensorCollapseShapeOp>::type;
+      std::is_same<TensorReshapeOpTy, tensor::CollapseShapeOp>::value,
+      tensor::ExpandShapeOp, tensor::CollapseShapeOp>::type;
   return b.create<ReverseReshapeOpTy>(reshapeOp.getLoc(),
                                       reshapeOp.getSrcType(), resultBuffer,
                                       reshapeOp.reassociation());
@@ -244,7 +244,7 @@
               }
               return nullptr;
             })
-            .Case<linalg::TensorCollapseShapeOp, linalg::TensorExpandShapeOp>(
+            .Case<tensor::CollapseShapeOp, tensor::ExpandShapeOp>(
                 [&](auto reshapeOp) {
                   return getReverseOfReshapeOp(b, reshapeOp, resultBuffer);
                 })
diff --git a/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp b/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
index bdd81a9..5f3a666 100644
--- a/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
+++ b/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
@@ -280,7 +280,7 @@
       reshapeOp.getSrcType(), {},
       resultBuffer.getType().cast<MemRefType>().getMemorySpace());
   using ReverseReshapeOpTy = typename std::conditional<
-      std::is_same<TensorReshapeOpTy, linalg::TensorCollapseShapeOp>::value,
+      std::is_same<TensorReshapeOpTy, tensor::CollapseShapeOp>::value,
       memref::ExpandShapeOp, memref::CollapseShapeOp>::type;
   return b.create<ReverseReshapeOpTy>(reshapeOp.getLoc(), memrefType,
                                       resultBuffer, reshapeOp.reassociation());
@@ -365,7 +365,7 @@
                   IREE::LinalgExt::LinalgExtOp, tensor::InsertSliceOp,
                   vector::TransferWriteOp>(
                 [&](auto op) { return resultBuffer; })
-            .Case<linalg::TensorCollapseShapeOp, linalg::TensorExpandShapeOp>(
+            .Case<tensor::CollapseShapeOp, tensor::ExpandShapeOp>(
                 [&](auto reshapeOp) {
                   return getReverseOfReshapeOp(b, reshapeOp, resultBuffer);
                 })
@@ -431,7 +431,7 @@
   auto reshapeResultType = getMemrefTypeForTensor(
       resultTensorType, {}, inputBufferType.getMemorySpace());
   using ReshapeOpTy = typename std::conditional<
-      std::is_same<TensorReshapeOpTy, linalg::TensorCollapseShapeOp>::value,
+      std::is_same<TensorReshapeOpTy, tensor::CollapseShapeOp>::value,
       memref::CollapseShapeOp, memref::ExpandShapeOp>::type;
   Value bufferReshape = b.create<ReshapeOpTy>(loc, reshapeResultType,
                                               inputBuffer, op.reassociation());
@@ -472,7 +472,7 @@
             tensor::CastOp>([&](auto singleResultOp) -> SmallVector<Value, 4> {
         return {getAliasingBufferForResult(b, singleResultOp, bvm)};
       })
-      .Case<linalg::TensorCollapseShapeOp, linalg::TensorExpandShapeOp>(
+      .Case<tensor::CollapseShapeOp, tensor::ExpandShapeOp>(
           [&](auto reshapeOp) -> SmallVector<Value, 4> {
             return {getAliasingBufferForReshapeResult(b, reshapeOp, bvm)};
           })
@@ -943,8 +943,8 @@
           }
           return convertScfIfOp(b, ifOp, bvm, plan);
         })
-        .Case<IREE::Flow::DispatchTensorLoadOp, linalg::TensorCollapseShapeOp,
-              linalg::TensorExpandShapeOp, tensor::ExtractSliceOp,
+        .Case<IREE::Flow::DispatchTensorLoadOp, tensor::CollapseShapeOp,
+              tensor::ExpandShapeOp, tensor::ExtractSliceOp,
               tensor::CastOp>([&](auto aliasingOp) {
           auto aliasingBuffers =
               getAliasingBuffersForResults(b, aliasingOp, bvm);
diff --git a/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp b/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp
index bb4dc4b..a1a5507 100644
--- a/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp
+++ b/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp
@@ -129,7 +129,7 @@
         Value one = b.create<arith::ConstantIndexOp>(loc, 1);
         std::array<Value, 3> returnValues = {one, one, one};
         for (auto ts : llvm::enumerate(currWorkloadPerWorkgroup)) {
-          returnValues[ts.index()] = linalg::applyMapToValues(
+          returnValues[ts.index()] = applyMapToValues(
               b, loc,
               AffineMap::get(0, 1,
                              b.getAffineSymbolExpr(0).ceilDiv(ts.value())),
diff --git a/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir b/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir
index 6db17a7..f9fe1a8 100644
--- a/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir
+++ b/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir
@@ -10,8 +10,8 @@
   %2 = hal.interface.binding.subspan @interface_io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x3x96xf32>
   // CHECK: %[[LOAD:.+]] = flow.dispatch.tensor.load %[[ARG]], {{.*}} : !flow.dispatch.tensor<readonly:3x3x96xf32> -> tensor<3x3x96xf32>
   %3 = flow.dispatch.tensor.load %1, offsets=[], sizes =[], strides=[] : !flow.dispatch.tensor<readonly:3x3x1x96xf32> -> tensor<3x3x1x96xf32>
-  %4 = linalg.tensor_collapse_shape %3 [[0, 1, 2, 3]] : tensor<3x3x1x96xf32> into tensor<864xf32>
-  %5 = linalg.tensor_expand_shape %4 [[0, 1, 2]] : tensor<864xf32> into tensor<3x3x96xf32>
+  %4 = tensor.collapse_shape %3 [[0, 1, 2, 3]] : tensor<3x3x1x96xf32> into tensor<864xf32>
+  %5 = tensor.expand_shape %4 [[0, 1, 2]] : tensor<864xf32> into tensor<3x3x96xf32>
   //  CHECK: flow.dispatch.tensor.store %[[LOAD]], {{.*}}
   flow.dispatch.tensor.store %5, %2, offsets = [%c0, %c0, %c0], sizes = [%c1, %c1, %c1], strides = [%c1, %c1, %c1] : tensor<3x3x96xf32> -> !flow.dispatch.tensor<writeonly:3x3x96xf32>
   return
@@ -34,10 +34,10 @@
   %1 = hal.interface.binding.subspan @interface_io::@arg0[%c0] : !flow.dispatch.tensor<readonly:6x3x1x96xf32>
   %2 = hal.interface.binding.subspan @interface_io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x3x96xf32>
   %3 = flow.dispatch.tensor.load %1, offsets = [%c3, %c0, %c0, %c0], sizes = [%c3, %c3, %c1, %c96], strides = [%c1, %c1, %c1, %c1] : !flow.dispatch.tensor<readonly:6x3x1x96xf32> -> tensor<3x3x1x96xf32>
-  // CHECK: linalg.tensor_collapse_shape
-  // CHECK: linalg.tensor_expand_shape
-  %4 = linalg.tensor_collapse_shape %3 [[0, 1, 2, 3]] : tensor<3x3x1x96xf32> into tensor<864xf32>
-  %5 = linalg.tensor_expand_shape %4 [[0, 1, 2]] : tensor<864xf32> into tensor<3x3x96xf32>
+  // CHECK: tensor.collapse_shape
+  // CHECK: tensor.expand_shape
+  %4 = tensor.collapse_shape %3 [[0, 1, 2, 3]] : tensor<3x3x1x96xf32> into tensor<864xf32>
+  %5 = tensor.expand_shape %4 [[0, 1, 2]] : tensor<864xf32> into tensor<3x3x96xf32>
   flow.dispatch.tensor.store %5, %2, offsets = [%c0, %c0, %c0], sizes = [%c1, %c1, %c1], strides = [%c1, %c1, %c1] : tensor<3x3x96xf32> -> !flow.dispatch.tensor<writeonly:3x3x96xf32>
   return
 }
@@ -59,10 +59,10 @@
   %1 = hal.interface.binding.subspan @interface_io::@arg0[%c0] : !flow.dispatch.tensor<readonly:?x?x96xf32>{%dim0, %dim1}
   %2 = hal.interface.binding.subspan @interface_io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:?x12x8xf32>{%dim2}
   %3 = flow.dispatch.tensor.load %1, offsets=[], sizes =[], strides=[] : !flow.dispatch.tensor<readonly:?x?x96xf32> -> tensor<?x?x96xf32>
-  // CHECK: linalg.tensor_collapse_shape
-  // CHECK: linalg.tensor_expand_shape
-  %4 = linalg.tensor_collapse_shape %3 [[0, 1], [2]] : tensor<?x?x96xf32> into tensor<?x96xf32>
-  %5 = linalg.tensor_expand_shape %4 [[0], [1, 2]] : tensor<?x96xf32> into tensor<?x12x8xf32>
+  // CHECK: tensor.collapse_shape
+  // CHECK: tensor.expand_shape
+  %4 = tensor.collapse_shape %3 [[0, 1], [2]] : tensor<?x?x96xf32> into tensor<?x96xf32>
+  %5 = tensor.expand_shape %4 [[0], [1, 2]] : tensor<?x96xf32> into tensor<?x12x8xf32>
   flow.dispatch.tensor.store %5, %2, offsets = [%c0, %c0, %c0], sizes = [%c1, %c1, %c1], strides = [%c1, %c1, %c1] : tensor<?x12x8xf32> -> !flow.dispatch.tensor<writeonly:?x12x8xf32>
   return
 }
diff --git a/iree/compiler/Codegen/Common/test/convert_to_destination_passing_style.mlir b/iree/compiler/Codegen/Common/test/convert_to_destination_passing_style.mlir
index 1712106..595a797 100644
--- a/iree/compiler/Codegen/Common/test/convert_to_destination_passing_style.mlir
+++ b/iree/compiler/Codegen/Common/test/convert_to_destination_passing_style.mlir
@@ -171,7 +171,7 @@
   %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:12xi32>
   %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32>
   %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
-  %3 = linalg.tensor_expand_shape %2 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
+  %3 = tensor.expand_shape %2 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
   flow.dispatch.tensor.store %3, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
   return
 }
@@ -183,7 +183,7 @@
 //  CHECK-DAG:   %[[ARG0:.+]] = hal.interface.binding.subspan @io::@arg0
 //  CHECK-DAG:   %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0
 //      CHECK:   %[[SOURCE:.+]] = flow.dispatch.tensor.load %[[ARG0]]
-//      CHECK:   %[[RESHAPE:.+]] = linalg.tensor_expand_shape %[[SOURCE]]
+//      CHECK:   %[[RESHAPE:.+]] = tensor.expand_shape %[[SOURCE]]
 //      CHECK:   flow.dispatch.tensor.store %[[RESHAPE]], %[[RET0]]
 
 // -----
@@ -197,7 +197,7 @@
   %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:12xi32>
   %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32>
   %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
-  %3 = linalg.tensor_expand_shape %2 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
+  %3 = tensor.expand_shape %2 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
   %4 = linalg.init_tensor [3, 4] : tensor<3x4xi32>
   %5 = linalg.generic {
     indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>],
@@ -219,7 +219,7 @@
 //  CHECK-DAG:   %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0
 //      CHECK:   %[[TARGET:.+]] = flow.dispatch.tensor.load %[[RET0]]
 //      CHECK:   %[[SOURCE:.+]] = flow.dispatch.tensor.load %[[ARG0]]
-//      CHECK:   %[[RESHAPE:.+]] = linalg.tensor_expand_shape %[[SOURCE]]
+//      CHECK:   %[[RESHAPE:.+]] = tensor.expand_shape %[[SOURCE]]
 //      CHECK:   %[[GENERIC:.+]] = linalg.generic
 // CHECK-SAME:       ins(%[[RESHAPE]]
 // CHECK-SAME:       outs(%[[TARGET]]
@@ -237,7 +237,7 @@
   %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32>
   %2 = hal.interface.binding.subspan @io::@ret1[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32>
   %3 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
-  %4 = linalg.tensor_expand_shape %3 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
+  %4 = tensor.expand_shape %3 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
   %5 = linalg.init_tensor [3, 4] : tensor<3x4xi32>
   %6 = linalg.generic {
     indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>],
@@ -262,7 +262,7 @@
 //  CHECK-DAG:   %[[TARGET:.+]] = flow.dispatch.tensor.load %[[RET0]]
 //  CHECK-DAG:   %[[RET1:.+]] = hal.interface.binding.subspan @io::@ret1
 //      CHECK:   %[[SOURCE:.+]] = flow.dispatch.tensor.load %[[ARG0]]
-//      CHECK:   %[[RESHAPE:.+]] = linalg.tensor_expand_shape %[[SOURCE]]
+//      CHECK:   %[[RESHAPE:.+]] = tensor.expand_shape %[[SOURCE]]
 //      CHECK:   %[[GENERIC:.+]] = linalg.generic
 // CHECK-SAME:       ins(%[[RESHAPE]]
 // CHECK-SAME:       outs(%[[TARGET]]
@@ -289,7 +289,7 @@
       %5 = arith.addi %arg0, %arg0 : i32
       linalg.yield %5 : i32
     } -> tensor<3x4xi32>
-  %5 = linalg.tensor_collapse_shape %4 [[0, 1]] : tensor<3x4xi32> into tensor<12xi32>
+  %5 = tensor.collapse_shape %4 [[0, 1]] : tensor<3x4xi32> into tensor<12xi32>
   flow.dispatch.tensor.store %5, %1, offsets = [], sizes = [], strides = [] : tensor<12xi32> -> !flow.dispatch.tensor<writeonly:12xi32>
   return
 }
@@ -302,11 +302,11 @@
 //  CHECK-DAG:   %[[RET0:.+]] = hal.interface.binding.subspan @io::@ret0
 //  CHECK-DAG:   %[[SOURCE:.+]] = flow.dispatch.tensor.load %[[ARG0]]
 //  CHECK-DAG:   %[[TARGET:.+]] = flow.dispatch.tensor.load %[[RET0]]
-//  CHECK-DAG:   %[[RESHAPE_EXPAND:.+]] = linalg.tensor_expand_shape %[[TARGET]] {{\[}}[0, 1]{{\]}}
+//  CHECK-DAG:   %[[RESHAPE_EXPAND:.+]] = tensor.expand_shape %[[TARGET]] {{\[}}[0, 1]{{\]}}
 //      CHECK:   %[[GENERIC:.+]] = linalg.generic
 // CHECK-SAME:       ins(%[[SOURCE]]
 // CHECK-SAME:       outs(%[[RESHAPE_EXPAND]]
-//      CHECK:   %[[RESHAPE_COLLAPSE:.+]] = linalg.tensor_collapse_shape %[[GENERIC]] {{\[}}[0, 1]{{\]}}
+//      CHECK:   %[[RESHAPE_COLLAPSE:.+]] = tensor.collapse_shape %[[GENERIC]] {{\[}}[0, 1]{{\]}}
 //      CHECK:   flow.dispatch.tensor.store %[[RESHAPE_COLLAPSE]], %[[RET0]]
 
 // -----
diff --git a/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
index 0ee94d6..e2484ab 100644
--- a/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
+++ b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
@@ -644,7 +644,7 @@
   %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:12xi32>
   %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32>
   %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
-  %3 = linalg.tensor_expand_shape %2 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
+  %3 = tensor.expand_shape %2 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
   flow.dispatch.tensor.store %3, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
   return
 }
@@ -669,7 +669,7 @@
   %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:12xi32>
   %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32>
   %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
-  %3 = linalg.tensor_expand_shape %2 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
+  %3 = tensor.expand_shape %2 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
   %4 = linalg.init_tensor [3, 4] : tensor<3x4xi32>
   %5 = linalg.generic {
     indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>],
@@ -707,7 +707,7 @@
   %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32>
   %2 = hal.interface.binding.subspan @io::@ret1[%c0] : !flow.dispatch.tensor<writeonly:3x4xi32>
   %3 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
-  %4 = linalg.tensor_expand_shape %3 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
+  %4 = tensor.expand_shape %3 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
   %5 = linalg.init_tensor [3, 4] : tensor<3x4xi32>
   %6 = linalg.generic {
     indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>],
@@ -757,7 +757,7 @@
       %5 = arith.addi %arg0, %arg0 : i32
       linalg.yield %5 : i32
     } -> tensor<3x4xi32>
-  %5 = linalg.tensor_collapse_shape %4 [[0, 1]] : tensor<3x4xi32> into tensor<12xi32>
+  %5 = tensor.collapse_shape %4 [[0, 1]] : tensor<3x4xi32> into tensor<12xi32>
   flow.dispatch.tensor.store %5, %1, offsets = [], sizes = [], strides = [] : tensor<12xi32> -> !flow.dispatch.tensor<writeonly:12xi32>
   return
 }
@@ -786,7 +786,7 @@
   %1 = hal.interface.binding.subspan @io::@arg1[%c0] : !flow.dispatch.tensor<readonly:2x3xf32>
   %2 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:1x3xf32>
   %3 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:1x1x2xf32> -> tensor<1x1x2xf32>
-  %4 = linalg.tensor_collapse_shape %3 [[0, 1], [2]] : tensor<1x1x2xf32> into tensor<1x2xf32>
+  %4 = tensor.collapse_shape %3 [[0, 1], [2]] : tensor<1x1x2xf32> into tensor<1x2xf32>
   %workgroup_size_x = hal.interface.workgroup.size[0] : index
   %workgroup_size_y = hal.interface.workgroup.size[1] : index
   %workgroup_id_x = hal.interface.workgroup.id[0] : index
@@ -1072,7 +1072,7 @@
   %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:1x5x3x1xf32>
   %1 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:5x5xf32>
   %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:1x5x3x1xf32> -> tensor<1x5x3x1xf32>
-  %3 = linalg.tensor_collapse_shape %2 [[0, 1], [2, 3]] : tensor<1x5x3x1xf32> into tensor<5x3xf32>
+  %3 = tensor.collapse_shape %2 [[0, 1], [2, 3]] : tensor<1x5x3x1xf32> into tensor<5x3xf32>
   %workgroup_size_x = hal.interface.workgroup.size[0] : index
   %workgroup_size_y = hal.interface.workgroup.size[1] : index
   %workgroup_id_x = hal.interface.workgroup.id[0] : index
@@ -1284,7 +1284,7 @@
   %0 = hal.interface.binding.subspan @io::@ro0[%c0] : !flow.dispatch.tensor<readonly:?x?xf32>{%dim0, %dim1}
   %1 = hal.interface.binding.subspan @io::@wo0[%c0] : !flow.dispatch.tensor<writeonly:?xf32>{%dim2}
   %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
-  %3 = linalg.tensor_collapse_shape %2 [[0, 1]]
+  %3 = tensor.collapse_shape %2 [[0, 1]]
       : tensor<?x?xf32> into tensor<?xf32>
   %4 = tensor.dim %3, %c0 : tensor<?xf32>
   %5 = linalg.init_tensor [%4] : tensor<?xf32>
@@ -1919,11 +1919,11 @@
         ^bb0(%arg3: f32, %arg4: f32):  // no predecessors
           linalg.yield %arg3 : f32
         } -> tensor<1x16x16x3x3x8xf32>
-        %19 = linalg.tensor_collapse_shape %18 [[0, 1, 2], [3, 4, 5]] : tensor<1x16x16x3x3x8xf32> into tensor<256x72xf32>
-        %20 = linalg.tensor_collapse_shape %14 [[0, 1, 2], [3]] : tensor<3x3x8x4xf32> into tensor<72x4xf32>
-        %21 = linalg.tensor_collapse_shape %16 [[0, 1, 2], [3]] : tensor<1x16x16x4xf32> into tensor<256x4xf32>
+        %19 = tensor.collapse_shape %18 [[0, 1, 2], [3, 4, 5]] : tensor<1x16x16x3x3x8xf32> into tensor<256x72xf32>
+        %20 = tensor.collapse_shape %14 [[0, 1, 2], [3]] : tensor<3x3x8x4xf32> into tensor<72x4xf32>
+        %21 = tensor.collapse_shape %16 [[0, 1, 2], [3]] : tensor<1x16x16x4xf32> into tensor<256x4xf32>
         %22 = linalg.matmul ins(%19, %20 : tensor<256x72xf32>, tensor<72x4xf32>) outs(%21 : tensor<256x4xf32>) -> tensor<256x4xf32>
-        %23 = linalg.tensor_expand_shape %22 [[0, 1, 2], [3]] : tensor<256x4xf32> into tensor<1x16x16x4xf32>
+        %23 = tensor.expand_shape %22 [[0, 1, 2], [3]] : tensor<256x4xf32> into tensor<1x16x16x4xf32>
         %24 = tensor.cast %23 : tensor<1x16x16x4xf32> to tensor<1x?x?x?xf32>
         flow.dispatch.tensor.store %24, %2, offsets = [0, %arg0, %arg1, %arg2], sizes = [1, %c16, %c16, %c4], strides = [1, 1, 1, 1] : tensor<1x?x?x?xf32> -> !flow.dispatch.tensor<writeonly:1x112x112x32xf32>
       }
diff --git a/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp b/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
index 8e5795f..9d1b8c1 100644
--- a/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
@@ -85,7 +85,8 @@
       arith::populateArithmeticToLLVMConversionPatterns(converter,
                                                         llvmPatterns);
       populateVectorToLLVMConversionPatterns(converter, llvmPatterns);
-      populateGpuToROCDLConversionPatterns(converter, llvmPatterns);
+      populateGpuToROCDLConversionPatterns(converter, llvmPatterns,
+                                           gpu::amd::Runtime::Unknown);
       LLVMConversionTarget target(getContext());
       populateStdToLLVMFuncOpConversionPattern(converter, llvmPatterns);
       configureGpuToROCDLConversionLegality(target);
diff --git a/iree/compiler/Dialect/Flow/Transforms/ConvertConv2D1x1ToMatmulPass.cpp b/iree/compiler/Dialect/Flow/Transforms/ConvertConv2D1x1ToMatmulPass.cpp
index ccf1568..74faa10 100644
--- a/iree/compiler/Dialect/Flow/Transforms/ConvertConv2D1x1ToMatmulPass.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/ConvertConv2D1x1ToMatmulPass.cpp
@@ -70,18 +70,18 @@
     Value output = convOp.getOutputOperand(0)->get();
     auto loc = convOp.getLoc();
 
-    Value reshapedInput = rewriter.create<linalg::TensorCollapseShapeOp>(
+    Value reshapedInput = rewriter.create<tensor::CollapseShapeOp>(
         loc, reshapedInputType, input, reassociationIndices);
-    Value reshapedFilter = rewriter.create<linalg::TensorCollapseShapeOp>(
+    Value reshapedFilter = rewriter.create<tensor::CollapseShapeOp>(
         loc, reshapedFilterType, filter, reassociationIndices);
-    Value reshapedOutput = rewriter.create<linalg::TensorCollapseShapeOp>(
+    Value reshapedOutput = rewriter.create<tensor::CollapseShapeOp>(
         loc, reshapedOutputType, output, reassociationIndices);
 
     auto matmulResult = rewriter.create<linalg::MatmulOp>(
         loc, reshapedOutputType, ArrayRef<Value>{reshapedInput, reshapedFilter},
         ArrayRef<Value>{reshapedOutput});
 
-    auto reshapedResult = rewriter.create<linalg::TensorExpandShapeOp>(
+    auto reshapedResult = rewriter.create<tensor::ExpandShapeOp>(
         loc, outputShapeType, matmulResult.getResults()[0],
         reassociationIndices);
 
diff --git a/iree/compiler/Dialect/Flow/Transforms/ConvertConv2DToImg2ColPass.cpp b/iree/compiler/Dialect/Flow/Transforms/ConvertConv2DToImg2ColPass.cpp
index 353bca7..9e7cff4 100644
--- a/iree/compiler/Dialect/Flow/Transforms/ConvertConv2DToImg2ColPass.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/ConvertConv2DToImg2ColPass.cpp
@@ -139,15 +139,14 @@
         RankedTensorType::get({outputShape[1] * outputShape[2], outputShape[3]},
                               outputShapeType.getElementType());
 
-    Value reshapedImg2ColTensor =
-        rewriter.create<linalg::TensorCollapseShapeOp>(
-            loc, reshapedImg2ColTensorType, img2ColTensor.getResult(0),
-            img2ColTensorReassociationIndices);
+    Value reshapedImg2ColTensor = rewriter.create<tensor::CollapseShapeOp>(
+        loc, reshapedImg2ColTensorType, img2ColTensor.getResult(0),
+        img2ColTensorReassociationIndices);
 
-    Value reshapedFilter = rewriter.create<linalg::TensorCollapseShapeOp>(
+    Value reshapedFilter = rewriter.create<tensor::CollapseShapeOp>(
         loc, reshapedFilterType, filter, filterAndOutputReassociationIndices);
 
-    Value reshapedOutput = rewriter.create<linalg::TensorCollapseShapeOp>(
+    Value reshapedOutput = rewriter.create<tensor::CollapseShapeOp>(
         loc, reshapedOutputType, output, filterAndOutputReassociationIndices);
 
     auto matmulResult = rewriter.create<linalg::MatmulOp>(
@@ -155,7 +154,7 @@
         ArrayRef<Value>{reshapedImg2ColTensor, reshapedFilter},
         ArrayRef<Value>{reshapedOutput});
 
-    auto reshapedResult = rewriter.create<linalg::TensorExpandShapeOp>(
+    auto reshapedResult = rewriter.create<tensor::ExpandShapeOp>(
         loc, outputShapeType, matmulResult.getResults()[0],
         filterAndOutputReassociationIndices);
 
@@ -308,14 +307,13 @@
          transposedOutputTensorShape[2] * transposedOutputTensorShape[3]},
         outputTensorType.getElementType());
 
-    Value reshapedImg2ColTensor =
-        rewriter.create<linalg::TensorCollapseShapeOp>(
-            loc, reshapedImg2ColTensorType, img2ColTensor.getResult(0),
-            img2ColTensorReassociationIndices);
-    Value reshapedFilterTensor = rewriter.create<linalg::TensorCollapseShapeOp>(
+    Value reshapedImg2ColTensor = rewriter.create<tensor::CollapseShapeOp>(
+        loc, reshapedImg2ColTensorType, img2ColTensor.getResult(0),
+        img2ColTensorReassociationIndices);
+    Value reshapedFilterTensor = rewriter.create<tensor::CollapseShapeOp>(
         loc, reshapedFilterTensorType, transposedFilter,
         filterReassociationIndice);
-    Value reshapedoutputTensor = rewriter.create<linalg::TensorCollapseShapeOp>(
+    Value reshapedoutputTensor = rewriter.create<tensor::CollapseShapeOp>(
         loc, reshapedOutputTensorType, transposedOutputTensor,
         outputReassociationIndice);
 
@@ -327,10 +325,9 @@
     SmallVector<ReassociationIndices> batchMatVecReassociationIndice = {{0, 1},
                                                                         {2, 3}};
 
-    Value batchMatVecResultReshaped =
-        rewriter.create<linalg::TensorExpandShapeOp>(
-            loc, transposedOutputTensor.getType(),
-            batchMatVecResult.getResult(0), batchMatVecReassociationIndice);
+    Value batchMatVecResultReshaped = rewriter.create<tensor::ExpandShapeOp>(
+        loc, transposedOutputTensor.getType(), batchMatVecResult.getResult(0),
+        batchMatVecReassociationIndice);
 
     auto transposedResult =
         transposeOperand(batchMatVecResultReshaped, {0, 2, 3, 1});
diff --git a/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgMatmulToMmt4D.cpp b/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgMatmulToMmt4D.cpp
index bf1e315..729937b 100644
--- a/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgMatmulToMmt4D.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgMatmulToMmt4D.cpp
@@ -44,7 +44,7 @@
       RankedTensorType::get(targetShape, inputType.getElementType());
   std::array<ReassociationIndices, 2> expandIndices = {
       ReassociationIndices{0, 1}, ReassociationIndices{2, 3}};
-  Value reshapedOperand = rewriter.create<linalg::TensorExpandShapeOp>(
+  Value reshapedOperand = rewriter.create<tensor::ExpandShapeOp>(
       loc, targetType, input, expandIndices);
   return reshapedOperand;
 }
@@ -103,7 +103,7 @@
       RankedTensorType::get(targetShape, inputType.getElementType());
   std::array<ReassociationIndices, 2> collapseIndices = {
       ReassociationIndices{0, 1}, ReassociationIndices{2, 3}};
-  Value reshapedOperand = rewriter.create<linalg::TensorCollapseShapeOp>(
+  Value reshapedOperand = rewriter.create<tensor::CollapseShapeOp>(
       loc, targetType, input, collapseIndices);
   return reshapedOperand;
 }
@@ -368,8 +368,9 @@
     // Canonicalization.
     {
       OwningRewritePatternList patterns(&getContext());
-      linalg::TensorExpandShapeOp::getCanonicalizationPatterns(patterns,
-                                                               context);
+      tensor::ExpandShapeOp::getCanonicalizationPatterns(patterns, context);
+      linalg::InitTensorOp::getCanonicalizationPatterns(patterns, context);
+      linalg::FillOp::getCanonicalizationPatterns(patterns, context);
       patterns.insert<FoldFillGenericOpPattern>(context);
       if (failed(applyPatternsAndFoldGreedily(getOperation(),
                                               std::move(patterns)))) {
diff --git a/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgTensorOps.cpp b/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgTensorOps.cpp
index 3e09e34..2bd32bf 100644
--- a/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgTensorOps.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/ConvertLinalgTensorOps.cpp
@@ -13,6 +13,7 @@
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/Dialect/Tensor/IR/Tensor.h"
 #include "mlir/IR/PatternMatch.h"
+#include "mlir/Interfaces/InferTypeOpInterface.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassManager.h"
 #include "mlir/Support/LLVM.h"
@@ -51,7 +52,10 @@
       return failure();
     }
     SmallVector<SmallVector<Value>> outputShape;
-    if (failed(reshapeOp.reifyResultShapes(rewriter, outputShape))) {
+    ReifyRankedShapedTypeOpInterface reifyShapedTypeInterface =
+        cast<ReifyRankedShapedTypeOpInterface>(reshapeOp.getOperation());
+    if (failed(reifyShapedTypeInterface.reifyResultShapes(rewriter,
+                                                          outputShape))) {
       return failure();
     }
     SmallVector<Value> outputDynamicShapes;
diff --git a/iree/compiler/Dialect/Flow/Transforms/ConvertToFlow.cpp b/iree/compiler/Dialect/Flow/Transforms/ConvertToFlow.cpp
index 993871f..156b9ed 100644
--- a/iree/compiler/Dialect/Flow/Transforms/ConvertToFlow.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/ConvertToFlow.cpp
@@ -52,7 +52,10 @@
       return failure();
     }
     SmallVector<SmallVector<Value>> outputShape;
-    if (failed(reshapeOp.reifyResultShapes(rewriter, outputShape))) {
+    ReifyRankedShapedTypeOpInterface reifyShapedTypeInterface =
+        cast<ReifyRankedShapedTypeOpInterface>(reshapeOp.getOperation());
+    if (failed(reifyShapedTypeInterface.reifyResultShapes(rewriter,
+                                                          outputShape))) {
       return failure();
     }
     SmallVector<Value> outputDynamicShapes;
@@ -105,10 +108,10 @@
     context->allowUnregisteredDialects(true);
     RewritePatternSet patterns(&getContext());
 
-    patterns.insert<
-        LinalgTensorReshapeToFlowTensorReshape<linalg::TensorCollapseShapeOp>,
-        LinalgTensorReshapeToFlowTensorReshape<linalg::TensorExpandShapeOp>>(
-        context);
+    patterns
+        .insert<LinalgTensorReshapeToFlowTensorReshape<tensor::CollapseShapeOp>,
+                LinalgTensorReshapeToFlowTensorReshape<tensor::ExpandShapeOp>>(
+            context);
     populateTensorToFlowPatternsBeforeDispatchFormation(context, patterns);
     IREE::Flow::TensorReshapeOp::getCanonicalizationPatterns(patterns, context);
 
diff --git a/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp b/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
index 6cd504c..a2f5564 100644
--- a/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
@@ -162,9 +162,9 @@
 /// Operations that are cloned into dispatch regions formed with other
 /// operations as roots.
 static bool isClonableIntoDispatchOp(Operation *op) {
-  if (isa<arith::IndexCastOp, linalg::InitTensorOp,
-          linalg::TensorCollapseShapeOp, linalg::TensorExpandShapeOp,
-          tensor::ExtractOp, tensor::ExtractSliceOp>(op)) {
+  if (isa<arith::IndexCastOp, linalg::InitTensorOp, tensor::CollapseShapeOp,
+          tensor::ExpandShapeOp, tensor::ExtractOp, tensor::ExtractSliceOp>(
+          op)) {
     return true;
   }
   if (auto constantOp = dyn_cast<arith::ConstantOp>(op)) {
diff --git a/iree/compiler/Dialect/Flow/Transforms/FusionOfTensorOps.cpp b/iree/compiler/Dialect/Flow/Transforms/FusionOfTensorOps.cpp
index f6f902b..9ac58a5 100644
--- a/iree/compiler/Dialect/Flow/Transforms/FusionOfTensorOps.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/FusionOfTensorOps.cpp
@@ -112,11 +112,10 @@
     // to the consumer linalg op.
     linalg::ControlElementwiseOpsFusionFn foldReshapeBetweenLinalgFn =
         [](const OpResult &producer, const OpOperand &consumer) {
-          auto collapseOp =
-              producer.getDefiningOp<linalg::TensorCollapseShapeOp>();
+          auto collapseOp = producer.getDefiningOp<tensor::CollapseShapeOp>();
           if (collapseOp)
             return collapseOp.src().getDefiningOp<LinalgOp>() != nullptr;
-          auto expandOp = producer.getDefiningOp<linalg::TensorExpandShapeOp>();
+          auto expandOp = producer.getDefiningOp<tensor::ExpandShapeOp>();
           if (expandOp)
             return expandOp.src().getDefiningOp<LinalgOp>() != nullptr;
           return false;
@@ -135,10 +134,14 @@
     OwningRewritePatternList reshapeCanonicalizations(&getContext());
     linalg::populateFoldUnitDimsReshapeOpsByLinearizationPatterns(
         reshapeCanonicalizations);
-    linalg::TensorCollapseShapeOp::getCanonicalizationPatterns(
+    tensor::CollapseShapeOp::getCanonicalizationPatterns(
         reshapeCanonicalizations, context);
-    linalg::TensorExpandShapeOp::getCanonicalizationPatterns(
-        reshapeCanonicalizations, context);
+    tensor::ExpandShapeOp::getCanonicalizationPatterns(reshapeCanonicalizations,
+                                                       context);
+    linalg::InitTensorOp::getCanonicalizationPatterns(reshapeCanonicalizations,
+                                                      context);
+    linalg::FillOp::getCanonicalizationPatterns(reshapeCanonicalizations,
+                                                context);
     if (failed(applyPatternsAndFoldGreedily(
             op->getRegions(), std::move(reshapeCanonicalizations)))) {
       return signalPassFailure();
@@ -147,10 +150,13 @@
     // Push the remaining reshapes down the graphs.
     OwningRewritePatternList pushReshapePatterns(&getContext());
     linalg::populatePushReshapeOpsPatterns(pushReshapePatterns);
-    linalg::TensorCollapseShapeOp::getCanonicalizationPatterns(
-        pushReshapePatterns, context);
-    linalg::TensorExpandShapeOp::getCanonicalizationPatterns(
-        pushReshapePatterns, context);
+    tensor::CollapseShapeOp::getCanonicalizationPatterns(pushReshapePatterns,
+                                                         context);
+    tensor::ExpandShapeOp::getCanonicalizationPatterns(pushReshapePatterns,
+                                                       context);
+    linalg::InitTensorOp::getCanonicalizationPatterns(pushReshapePatterns,
+                                                      context);
+    linalg::FillOp::getCanonicalizationPatterns(pushReshapePatterns, context);
     if (failed(applyPatternsAndFoldGreedily(op->getRegions(),
                                             std::move(pushReshapePatterns)))) {
       return signalPassFailure();
diff --git a/iree/compiler/Dialect/Flow/Transforms/PadTensorToSubTensorInsert.cpp b/iree/compiler/Dialect/Flow/Transforms/PadTensorToSubTensorInsert.cpp
index 9667be2..7d28c51 100644
--- a/iree/compiler/Dialect/Flow/Transforms/PadTensorToSubTensorInsert.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/PadTensorToSubTensorInsert.cpp
@@ -80,7 +80,7 @@
       };
       expr = addValueOrAttr(expr, lowPad[dim]);
       expr = addValueOrAttr(expr, highPad[dim]);
-      Value v = linalg::applyMapToValues(
+      Value v = applyMapToValues(
           rewriter, loc, AffineMap::get(1, numSymbols, expr), mapValues)[0];
       if (auto cst = v.getDefiningOp<arith::ConstantOp>()) {
         outputShape.push_back(cst.getValue());
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/conv1x1_to_matmul.mlir b/iree/compiler/Dialect/Flow/Transforms/test/conv1x1_to_matmul.mlir
index 72312f2..fc26e2b 100644
--- a/iree/compiler/Dialect/Flow/Transforms/test/conv1x1_to_matmul.mlir
+++ b/iree/compiler/Dialect/Flow/Transforms/test/conv1x1_to_matmul.mlir
@@ -12,9 +12,9 @@
 // CHECK: %[[INPUT:.+]]: tensor<1x4x5x2xf32>
 // CHECK: %[[FILTER:.+]]: tensor<1x1x2x7xf32>
 // CHECK: %[[OTUPUT:.+]] = linalg.init_tensor [1, 4, 5, 7] : tensor<1x4x5x7xf32>
-// CHECK: %[[RESHAPED_INPUT:.+]] = linalg.tensor_collapse_shape %[[INPUT]] {{\[}}[0, 1, 2], [3]] : tensor<1x4x5x2xf32> into tensor<20x2xf32>
-// CHECK: %[[RESHAPED_FILTER:.+]] = linalg.tensor_collapse_shape %[[FILTER]] {{\[}}[0, 1, 2], [3]] : tensor<1x1x2x7xf32> into tensor<2x7xf32>
-// CHECK: %[[RESHAPED_OUTPUT:.+]] = linalg.tensor_collapse_shape %[[OTUPUT]] {{\[}}[0, 1, 2], [3]] : tensor<1x4x5x7xf32> into tensor<20x7xf32>
+// CHECK: %[[RESHAPED_INPUT:.+]] = tensor.collapse_shape %[[INPUT]] {{\[}}[0, 1, 2], [3]] : tensor<1x4x5x2xf32> into tensor<20x2xf32>
+// CHECK: %[[RESHAPED_FILTER:.+]] = tensor.collapse_shape %[[FILTER]] {{\[}}[0, 1, 2], [3]] : tensor<1x1x2x7xf32> into tensor<2x7xf32>
+// CHECK: %[[RESHAPED_OUTPUT:.+]] = tensor.collapse_shape %[[OTUPUT]] {{\[}}[0, 1, 2], [3]] : tensor<1x4x5x7xf32> into tensor<20x7xf32>
 // CHECK: %[[MATMUL_RESULT:.+]] = linalg.matmul ins(%[[RESHAPED_INPUT]], %[[RESHAPED_FILTER]] : tensor<20x2xf32>, tensor<2x7xf32>) outs(%[[RESHAPED_OUTPUT]] : tensor<20x7xf32>)
-// CHECK: %[[RESULT:.+]] = linalg.tensor_expand_shape %[[MATMUL_RESULT]] {{\[}}[0, 1, 2], [3]] : tensor<20x7xf32> into tensor<1x4x5x7xf32>
+// CHECK: %[[RESULT:.+]] = tensor.expand_shape %[[MATMUL_RESULT]] {{\[}}[0, 1, 2], [3]] : tensor<20x7xf32> into tensor<1x4x5x7xf32>
 // CHECK: return %[[RESULT]]
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/conv2d_to_img2col.mlir b/iree/compiler/Dialect/Flow/Transforms/test/conv2d_to_img2col.mlir
index 268e87a..3c86ba8 100644
--- a/iree/compiler/Dialect/Flow/Transforms/test/conv2d_to_img2col.mlir
+++ b/iree/compiler/Dialect/Flow/Transforms/test/conv2d_to_img2col.mlir
@@ -19,16 +19,16 @@
 //           CHECK-SAME: #[[MAP1]]
 //                CHECK: ^bb0(%[[IN_DATA:.+]]: f32, %[[OUT_DATA:.+]]: f32)
 //                CHECK: linalg.yield %[[IN_DATA]] : f32
-//      CHECK-DAG: %[[RESHAPED_INIT_COL_TENSOR:.+]] = linalg.tensor_collapse_shape %[[COL_TENSOR]]
+//      CHECK-DAG: %[[RESHAPED_INIT_COL_TENSOR:.+]] = tensor.collapse_shape %[[COL_TENSOR]]
 //           CHECK-SAME: [0, 1, 2], [3, 4, 5]
 //           CHECK-SAME: tensor<1x14x14x3x3x4xf32> into tensor<196x36xf32>
-//      CHECK-DAG: %[[RESHAPED_FILTER:.+]] = linalg.tensor_collapse_shape %[[FILTER]]
+//      CHECK-DAG: %[[RESHAPED_FILTER:.+]] = tensor.collapse_shape %[[FILTER]]
 //           CHECK-SAME: [0, 1, 2], [3]
 //           CHECK-SAME: tensor<3x3x4x16xf32> into tensor<36x16xf32>
-//      CHECK-DAG: %[[RESHAPED_OUTPUT:.+]] = linalg.tensor_collapse_shape %[[OUTPUT]]
+//      CHECK-DAG: %[[RESHAPED_OUTPUT:.+]] = tensor.collapse_shape %[[OUTPUT]]
 //           CHECK-SAME: [0, 1, 2], [3]
 //      CHECK: %[[MATMUL_RESULT:.+]] = linalg.matmul ins(%[[RESHAPED_INIT_COL_TENSOR]], %[[RESHAPED_FILTER]] : tensor<196x36xf32>, tensor<36x16xf32>) outs(%[[RESHAPED_OUTPUT]] : tensor<196x16xf32>)
-//      CHECK: %[[RESULT:.+]] = linalg.tensor_expand_shape %[[MATMUL_RESULT]] {{\[}}[0, 1, 2], [3]] : tensor<196x16xf32> into tensor<1x14x14x16xf32>
+//      CHECK: %[[RESULT:.+]] = tensor.expand_shape %[[MATMUL_RESULT]] {{\[}}[0, 1, 2], [3]] : tensor<196x16xf32> into tensor<1x14x14x16xf32>
 //      CHECK: return %[[RESULT]]
 
 // -----
@@ -84,14 +84,14 @@
 // CHECK-NEXT:      ^bb0(%{{.*}}: f32, %{{.*}}: f32):
 // CHECK-NEXT:         linalg.yield
 // CHECK-NEXT:    } -> tensor<1x16x112x112x3x3xf32>
-//      CHECK: %[[COL_TENSOR_R:.+]] = linalg.tensor_collapse_shape %[[COL_TENSOR]]
+//      CHECK: %[[COL_TENSOR_R:.+]] = tensor.collapse_shape %[[COL_TENSOR]]
 // CHECK-SAME:    tensor<1x16x112x112x3x3xf32> into tensor<16x12544x9xf32>
-//      CHECK: %[[FILTER_T_R:.+]] = linalg.tensor_collapse_shape %[[FILTER_T]]
+//      CHECK: %[[FILTER_T_R:.+]] = tensor.collapse_shape %[[FILTER_T]]
 // CHECK-SAME:    tensor<16x3x3xf32> into tensor<16x9xf32>
-//      CHECK: %[[OUTPUT_T_R:.+]] = linalg.tensor_collapse_shape %[[OUTPUT_T]]
+//      CHECK: %[[OUTPUT_T_R:.+]] = tensor.collapse_shape %[[OUTPUT_T]]
 // CHECK-SAME:    tensor<1x16x112x112xf32> into tensor<16x12544xf32>
 //      CHECK: %[[BMV_RESULT:.+]] = linalg.batch_matvec ins(%[[COL_TENSOR_R]], %[[FILTER_T_R]] : tensor<16x12544x9xf32>, tensor<16x9xf32>) outs(%[[OUTPUT_T_R]] : tensor<16x12544xf32>) -> tensor<16x12544xf32>
-//      CHECK: %[[RESULT_R:.+]] = linalg.tensor_expand_shape %[[BMV_RESULT]]
+//      CHECK: %[[RESULT_R:.+]] = tensor.expand_shape %[[BMV_RESULT]]
 // CHECK-SAME:    tensor<16x12544xf32> into tensor<1x16x112x112xf32>
 //      CHECK: %[[RESULT_INIT:.+]] = linalg.init_tensor [1, 112, 112, 16] : tensor<1x112x112x16xf32>
 //      CHECK: %[[RESULT:.+]] = linalg.generic
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/convert_linalg_tensor_ops_before.mlir b/iree/compiler/Dialect/Flow/Transforms/test/convert_linalg_tensor_ops_before.mlir
index 1fe2725..957bbb8 100644
--- a/iree/compiler/Dialect/Flow/Transforms/test/convert_linalg_tensor_ops_before.mlir
+++ b/iree/compiler/Dialect/Flow/Transforms/test/convert_linalg_tensor_ops_before.mlir
@@ -3,9 +3,9 @@
 func @tensor_reshape(%arg0 : tensor<?x4x?x5x?x6xf32>, %arg1 : tensor<20x?x40xf32>)
     -> (tensor<?x5x?xf32>, tensor<5x4x?x4x2x4x5xf32>)
 {
-  %0 = linalg.tensor_collapse_shape %arg0 [[0, 1, 2], [3], [4, 5]]
+  %0 = tensor.collapse_shape %arg0 [[0, 1, 2], [3], [4, 5]]
       : tensor<?x4x?x5x?x6xf32> into tensor<?x5x?xf32>
-  %1 = linalg.tensor_expand_shape %arg1 [[0, 1], [2, 3], [4, 5, 6]]
+  %1 = tensor.expand_shape %arg1 [[0, 1], [2, 3], [4, 5, 6]]
       : tensor<20x?x40xf32> into tensor<5x4x?x4x2x4x5xf32>
   return %0, %1 : tensor<?x5x?xf32>, tensor<5x4x?x4x2x4x5xf32>
 }
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir
index fa7e415..5ddf2a7 100644
--- a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir
+++ b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir
@@ -265,7 +265,7 @@
   %cst = arith.constant 0.0 : f32
   %c0 = arith.constant 0 : index
   %c1 = arith.constant 1 : index
-  %0 = linalg.tensor_expand_shape %lhs [[0, 1]]
+  %0 = tensor.expand_shape %lhs [[0, 1]]
     : tensor<?xf32> into tensor<?x4xf32>
   %m = tensor.dim %0, %c0 : tensor<?x4xf32>
   %n1 = tensor.dim %rhs1, %c1 : tensor<4x?xf32>
@@ -543,14 +543,14 @@
 func @inline_dag_1(
     %arg0: tensor<?xf32>, %arg1 : tensor<1x?xf32>, %arg2 : tensor<i32>,
     %arg3 : index) -> tensor<?xf32> {
-  %0 = linalg.tensor_expand_shape %arg0 [[0, 1]] : tensor<?xf32> into tensor<1x?xf32>
+  %0 = tensor.expand_shape %arg0 [[0, 1]] : tensor<?xf32> into tensor<1x?xf32>
   %1 = tensor.extract_slice %0[0, 20] [1, %arg3] [1, 1] : tensor<1x?xf32> to tensor<1x?xf32>
-  %2 = linalg.tensor_collapse_shape %1 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
-  %3 = linalg.tensor_collapse_shape %arg1 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
+  %2 = tensor.collapse_shape %1 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
+  %3 = tensor.collapse_shape %arg1 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
   %4 = tensor.extract_slice %0[0, 10] [1, %arg3] [1, 1] : tensor<1x?xf32> to tensor<1x?xf32>
-  %5 = linalg.tensor_collapse_shape %4 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
+  %5 = tensor.collapse_shape %4 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
   %6 = tensor.extract_slice %0[0, 0] [1, %arg3] [1, 1] : tensor<1x?xf32> to tensor<1x?xf32>
-  %7 = linalg.tensor_collapse_shape %6 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
+  %7 = tensor.collapse_shape %6 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
   %8 = linalg.init_tensor [%arg3] : tensor<?xf32>
   %9 = linalg.generic {
       indexing_maps = [#map0, #map1, #map1, #map1, #map1, #map1],
@@ -580,14 +580,14 @@
 //       CHECK:     %[[LEAF1:.+]] = flow.dispatch.tensor.load %[[ARG4]], {{.*}}
 //       CHECK:     %[[LEAF2:.+]] = flow.dispatch.tensor.load %[[ARG5]], {{.*}}
 //       CHECK:     %[[LEAF3:.+]] = flow.dispatch.tensor.load %[[ARG6]], {{.*}}
-//       CHECK:     %[[OP1:.+]] = linalg.tensor_expand_shape %[[LEAF2]]
-//       CHECK:     %[[OP2:.+]] = linalg.tensor_collapse_shape %[[LEAF1]]
+//       CHECK:     %[[OP1:.+]] = tensor.expand_shape %[[LEAF2]]
+//       CHECK:     %[[OP2:.+]] = tensor.collapse_shape %[[LEAF1]]
 //       CHECK:     %[[OP3:.+]] = tensor.extract_slice %[[OP1]][0, 0]
 //       CHECK:     %[[OP4:.+]] = tensor.extract_slice %[[OP1]][0, 10]
 //       CHECK:     %[[OP5:.+]] = tensor.extract_slice %[[OP1]][0, 20]
-//       CHECK:     %[[OP6:.+]] = linalg.tensor_collapse_shape %[[OP3]]
-//       CHECK:     %[[OP7:.+]] = linalg.tensor_collapse_shape %[[OP4]]
-//       CHECK:     %[[OP8:.+]] = linalg.tensor_collapse_shape %[[OP5]]
+//       CHECK:     %[[OP6:.+]] = tensor.collapse_shape %[[OP3]]
+//       CHECK:     %[[OP7:.+]] = tensor.collapse_shape %[[OP4]]
+//       CHECK:     %[[OP8:.+]] = tensor.collapse_shape %[[OP5]]
 
 // -----
 
@@ -596,16 +596,16 @@
 func @inline_dag_2(
     %arg0: tensor<?xf32>, %arg1 : tensor<1x?xf32>, %arg2 : tensor<i32>,
     %arg3 : index) -> tensor<?xf32> {
-  %0 = linalg.tensor_expand_shape %arg0 [[0, 1]] : tensor<?xf32> into tensor<1x?xf32>
+  %0 = tensor.expand_shape %arg0 [[0, 1]] : tensor<?xf32> into tensor<1x?xf32>
   %1 = tensor.extract_slice %0[0, 20] [1, %arg3] [1, 1] : tensor<1x?xf32> to tensor<1x?xf32>
-  %2 = linalg.tensor_collapse_shape %arg1 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
+  %2 = tensor.collapse_shape %arg1 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
   br ^bb1
 ^bb1:
-  %3 = linalg.tensor_collapse_shape %1 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
+  %3 = tensor.collapse_shape %1 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
   %4 = tensor.extract_slice %0[0, 10] [1, %arg3] [1, 1] : tensor<1x?xf32> to tensor<1x?xf32>
-  %5 = linalg.tensor_collapse_shape %4 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
+  %5 = tensor.collapse_shape %4 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
   %6 = tensor.extract_slice %0[0, 0] [1, %arg3] [1, 1] : tensor<1x?xf32> to tensor<1x?xf32>
-  %7 = linalg.tensor_collapse_shape %6 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
+  %7 = tensor.collapse_shape %6 [[0, 1]] : tensor<1x?xf32> into tensor<?xf32>
   %8 = linalg.init_tensor [%arg3] : tensor<?xf32>
   %9 = linalg.generic {
       indexing_maps = [#map0, #map1, #map1, #map1, #map1, #map1],
@@ -635,14 +635,14 @@
 //       CHECK:     %[[LEAF1:.+]] = flow.dispatch.tensor.load %[[ARG4]], {{.*}}
 //       CHECK:     %[[LEAF2:.+]] = flow.dispatch.tensor.load %[[ARG5]], {{.*}}
 //       CHECK:     %[[LEAF3:.+]] = flow.dispatch.tensor.load %[[ARG6]], {{.*}}
-//       CHECK:     %[[OP1:.+]] = linalg.tensor_expand_shape %[[LEAF2]]
-//       CHECK:     %[[OP2:.+]] = linalg.tensor_collapse_shape %[[LEAF1]]
+//       CHECK:     %[[OP1:.+]] = tensor.expand_shape %[[LEAF2]]
+//       CHECK:     %[[OP2:.+]] = tensor.collapse_shape %[[LEAF1]]
 //       CHECK:     %[[OP3:.+]] = tensor.extract_slice %[[OP1]][0, 0]
 //       CHECK:     %[[OP4:.+]] = tensor.extract_slice %[[OP1]][0, 10]
 //       CHECK:     %[[OP5:.+]] = tensor.extract_slice %[[OP1]][0, 20]
-//       CHECK:     %[[OP6:.+]] = linalg.tensor_collapse_shape %[[OP3]]
-//       CHECK:     %[[OP7:.+]] = linalg.tensor_collapse_shape %[[OP4]]
-//       CHECK:     %[[OP8:.+]] = linalg.tensor_collapse_shape %[[OP5]]
+//       CHECK:     %[[OP6:.+]] = tensor.collapse_shape %[[OP3]]
+//       CHECK:     %[[OP7:.+]] = tensor.collapse_shape %[[OP4]]
+//       CHECK:     %[[OP8:.+]] = tensor.collapse_shape %[[OP5]]
 
 // -----
 
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/matmul_to_mmt4d.mlir b/iree/compiler/Dialect/Flow/Transforms/test/matmul_to_mmt4d.mlir
index 1001bb0..19caa34 100644
--- a/iree/compiler/Dialect/Flow/Transforms/test/matmul_to_mmt4d.mlir
+++ b/iree/compiler/Dialect/Flow/Transforms/test/matmul_to_mmt4d.mlir
@@ -8,11 +8,11 @@
 // CHECK-DAG:#[[MAP1:.+]] = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>
 // CHECK-DAG:#[[MAP2:.+]] = affine_map<(d0, d1, d2, d3) -> (d1, d3, d0, d2)>
 //      CHECK: @check_mmt4d_f32_static_nopad(%[[LHS:.+]]: tensor<24x8xf32>, %[[RHS:.+]]: tensor<8x32xf32>, %[[DST:.+]]: tensor<24x32xf32>)
-//      CHECK: %[[LHS4D:.+]] = linalg.tensor_expand_shape %[[LHS]]
+//      CHECK: %[[LHS4D:.+]] = tensor.expand_shape %[[LHS]]
 // CHECK-SAME:   tensor<24x8xf32> into tensor<3x8x4x2xf32>
-//      CHECK: %[[RHS4D:.+]] = linalg.tensor_expand_shape %[[RHS]]
+//      CHECK: %[[RHS4D:.+]] = tensor.expand_shape %[[RHS]]
 // CHECK-SAME:   tensor<8x32xf32> into tensor<4x2x8x4xf32>
-//      CHECK: %[[DST4D:.+]] = linalg.tensor_expand_shape %[[DST]]
+//      CHECK: %[[DST4D:.+]] = tensor.expand_shape %[[DST]]
 // CHECK-SAME:   tensor<24x32xf32> into tensor<3x8x8x4xf32>
 //      CHECK: %[[LHS4DT_INIT:.+]] = linalg.init_tensor [3, 4, 8, 2] : tensor<3x4x8x2xf32>
 //      CHECK: %[[LHS4DT:.+]] = linalg.generic
@@ -47,7 +47,7 @@
 // CHECK-NEXT:    ^bb0(%{{.*}}: f32, %{{.*}}: f32):
 // CHECK-NEXT:           linalg.yield %arg3 : f32
 // CHECK-NEXT:    } -> tensor<3x8x8x4xf32>
-//      CHECK: %[[RESULT:.+]] = linalg.tensor_collapse_shape %[[MMT4DT]]
+//      CHECK: %[[RESULT:.+]] = tensor.collapse_shape %[[MMT4DT]]
 // CHECK-SAME:    tensor<3x8x8x4xf32> into tensor<24x32xf32>
 //      CHECK: return %[[RESULT]] : tensor<24x32xf32>
 
@@ -64,9 +64,9 @@
 // CHECK-DAG:#[[MAP2:.+]] = affine_map<(d0, d1, d2, d3) -> (d1, d3, d0, d2)>
 //      CHECK: @check_mmt4d_with_init_tensor_and_fill(%[[LHS:.+]]: tensor<24x8xf32>, %[[RHS:.+]]: tensor<8x32xf32>)
 //      CHECK: %[[ZERO:.+]] = arith.constant 0.000000e+00 : f32
-//      CHECK: %[[LHS4D:.+]] = linalg.tensor_expand_shape %[[LHS]]
+//      CHECK: %[[LHS4D:.+]] = tensor.expand_shape %[[LHS]]
 // CHECK-SAME:   tensor<24x8xf32> into tensor<3x8x4x2xf32>
-//      CHECK: %[[RHS4D:.+]] = linalg.tensor_expand_shape %[[RHS]]
+//      CHECK: %[[RHS4D:.+]] = tensor.expand_shape %[[RHS]]
 // CHECK-SAME:   tensor<8x32xf32> into tensor<4x2x8x4xf32>
 //      CHECK: %[[DST_INIT:.+]] = linalg.init_tensor [3, 8, 8, 4] : tensor<3x8x8x4xf32>
 //      CHECK: [[DST:.+]] linalg.fill(%[[ZERO:.+]], %[[DST_INIT]])
@@ -84,15 +84,15 @@
 //      CHECK: tensor<5x2xi8> to tensor<6x4xi8>
 //      CHECK: %[[ACCPAD:.+]] = linalg.pad_tensor %[[ACC]] low[0, 0] high[5, 2]
 //      CHECK: tensor<3x2xi32> to tensor<8x4xi32>
-//      CHECK: %[[LHS4D:.+]] = linalg.tensor_expand_shape %[[LHSPAD]]
+//      CHECK: %[[LHS4D:.+]] = tensor.expand_shape %[[LHSPAD]]
 // CHECK-SAME: tensor<8x6xi8> into tensor<1x8x3x2xi8>
-//      CHECK: %[[RHS4D:.+]] = linalg.tensor_expand_shape %[[RHSPAD]]
+//      CHECK: %[[RHS4D:.+]] = tensor.expand_shape %[[RHSPAD]]
 // CHECK-SAME: tensor<6x4xi8> into tensor<3x2x1x4xi8>
-//      CHECK: %[[ACC4D:.+]] = linalg.tensor_expand_shape %[[ACCPAD]]
+//      CHECK: %[[ACC4D:.+]] = tensor.expand_shape %[[ACCPAD]]
 // CHECK-SAME: tensor<8x4xi32> into tensor<1x8x1x4xi32>
 //  ... After the above padding, we are reduced to the same stuff as we have
 //  ... already checked in the above testcases, so we skip checking that again.
-//      CHECK: %[[RESPAD:.+]] = linalg.tensor_collapse_shape
+//      CHECK: %[[RESPAD:.+]] = tensor.collapse_shape
 // CHECK-SAME: tensor<1x8x1x4xi32> into tensor<8x4xi32>
 //      CHECK: %[[RES:.+]] = tensor.extract_slice %[[RESPAD]][0, 0] [3, 2] [1, 1]
 // CHECK-SAME: tensor<8x4xi32> to tensor<3x2xi32>
@@ -113,15 +113,15 @@
 //      CHECK: tensor<?x?xi8> to tensor<?x?xi8>
 //      CHECK: %[[ACCPAD:.+]] = linalg.pad_tensor %[[ACC]] low[0, 0] high[
 //      CHECK: tensor<?x?xi32> to tensor<?x?xi32>
-//      CHECK: %[[LHS4D:.+]] = linalg.tensor_expand_shape %[[LHSPAD]]
+//      CHECK: %[[LHS4D:.+]] = tensor.expand_shape %[[LHSPAD]]
 // CHECK-SAME: tensor<?x?xi8> into tensor<?x8x?x2xi8>
-//      CHECK: %[[RHS4D:.+]] = linalg.tensor_expand_shape %[[RHSPAD]]
+//      CHECK: %[[RHS4D:.+]] = tensor.expand_shape %[[RHSPAD]]
 // CHECK-SAME: tensor<?x?xi8> into tensor<?x2x?x4xi8>
-//      CHECK: %[[ACC4D:.+]] = linalg.tensor_expand_shape %[[ACCPAD]]
+//      CHECK: %[[ACC4D:.+]] = tensor.expand_shape %[[ACCPAD]]
 // CHECK-SAME: tensor<?x?xi32> into tensor<?x8x?x4xi32>
 //  ... After the above padding, we are reduced to the same stuff as we have
 //  ... already checked in the above testcases, so we skip checking that again.
-//      CHECK: %[[RESPAD:.+]] = linalg.tensor_collapse_shape
+//      CHECK: %[[RESPAD:.+]] = tensor.collapse_shape
 // CHECK-SAME: tensor<?x8x?x4xi32> into tensor<?x?xi32>
 //      CHECK: %[[RES:.+]] = tensor.extract_slice %[[RESPAD]][0, 0] [{{.*}}] [1, 1]
 // CHECK-SAME: tensor<?x?xi32> to tensor<?x?xi32>
diff --git a/iree/compiler/InputConversion/MHLO/ConvertMHLOToLinalgExt.cpp b/iree/compiler/InputConversion/MHLO/ConvertMHLOToLinalgExt.cpp
index df0ff46..511e3ff 100644
--- a/iree/compiler/InputConversion/MHLO/ConvertMHLOToLinalgExt.cpp
+++ b/iree/compiler/InputConversion/MHLO/ConvertMHLOToLinalgExt.cpp
@@ -257,7 +257,7 @@
     map.emplace_back(1, indicesRank - 1);
     auto resultType = RankedTensorType::get({batchSize, shape.back()},
                                             indicesType.getElementType());
-    indices = b.create<linalg::TensorCollapseShapeOp>(resultType, indices, map);
+    indices = b.create<tensor::CollapseShapeOp>(resultType, indices, map);
 
     auto updateShape = updatesType.getShape().drop_front(shape.size() - 1);
     SmallVector<int64_t> collapsedUpdateShape = {batchSize};
@@ -269,7 +269,7 @@
     for (auto i : llvm::seq<int64_t>(indicesRank - 1, updatesType.getRank())) {
       map.emplace_back(1, i);
     }
-    updates = b.create<linalg::TensorCollapseShapeOp>(resultType, updates, map);
+    updates = b.create<tensor::CollapseShapeOp>(resultType, updates, map);
 
     return success();
   }
diff --git a/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp b/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp
index b6bec24..fa18aa4 100644
--- a/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp
+++ b/iree/compiler/InputConversion/MHLO/LegalizeInputTypes.cpp
@@ -81,8 +81,8 @@
                                       BlockAndValueMapping &mapping,
                                       OpBuilder &builder) {
   if (isa<linalg::LinalgDialect>(oldOp->getDialect()) &&
-      !isa<linalg::TensorCollapseShapeOp>(oldOp) &&
-      !isa<linalg::TensorExpandShapeOp>(oldOp)) {
+      !isa<tensor::CollapseShapeOp>(oldOp) &&
+      !isa<tensor::ExpandShapeOp>(oldOp)) {
     // Currently we assume all Linalg structured ops only contain valid types.
     builder.clone(*oldOp, mapping);
     return success();
diff --git a/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir b/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir
index 10115da..bba02f6 100644
--- a/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir
+++ b/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir
@@ -336,9 +336,9 @@
 // CHECK:         %[[ARG0:[a-zA-Z0-9]+]]
 // CHECK:         %[[ARG1:[a-zA-Z0-9]+]]
 // CHECK:         %[[ARG2:[a-zA-Z0-9]+]]
-// CHECK:         %[[COLLAPSED_INDICES:.+]] = linalg.tensor_collapse_shape
+// CHECK:         %[[COLLAPSED_INDICES:.+]] = tensor.collapse_shape
 // CHECK-SAME:        %[[ARG1]] {{\[}}[0, 1], [2]] : tensor<3x4x1xi32> into tensor<12x1xi32>
-// CHECK:         %[[COLLAPSED_UPDATES:.+]] = linalg.tensor_collapse_shape
+// CHECK:         %[[COLLAPSED_UPDATES:.+]] = tensor.collapse_shape
 // CHECK-SAME:        %[[ARG2]] {{\[}}[0, 1]] : tensor<3x4xi32> into tensor<12xi32>
 // CHECK:         %[[SCATTER:.+]] = iree_linalg_ext.scatter
 // CHECK-SAME:       ins(%[[COLLAPSED_UPDATES]], %[[COLLAPSED_INDICES]] : tensor<12xi32>, tensor<12x1xi32>)
@@ -369,9 +369,9 @@
 // CHECK:         %[[ARG0:[a-zA-Z0-9]+]]
 // CHECK:         %[[ARG1:[a-zA-Z0-9]+]]
 // CHECK:         %[[ARG2:[a-zA-Z0-9]+]]
-// CHECK:         %[[COLLAPSED_INDICES:.+]] = linalg.tensor_collapse_shape
+// CHECK:         %[[COLLAPSED_INDICES:.+]] = tensor.collapse_shape
 // CHECK-SAME:        %[[ARG1]] {{\[}}[0, 1], [2]] : tensor<?x3x2xi32> into tensor<?x2xi32>
-// CHECK:         %[[COLLAPSED_UPDATES:.+]] = linalg.tensor_collapse_shape
+// CHECK:         %[[COLLAPSED_UPDATES:.+]] = tensor.collapse_shape
 // CHECK-SAME:        %[[ARG2]] {{\[}}[0, 1], [2]] : tensor<?x3x512xi32> into tensor<?x512xi32>
 // CHECK:         %[[SCATTER:.+]] = iree_linalg_ext.scatter
 // CHECK-SAME:        ins(%[[COLLAPSED_UPDATES]], %[[COLLAPSED_INDICES]] : tensor<?x512xi32>, tensor<?x2xi32>)
diff --git a/iree/compiler/InputConversion/MHLO/test/legalize_input_types.mlir b/iree/compiler/InputConversion/MHLO/test/legalize_input_types.mlir
index 313df38..2467f56 100644
--- a/iree/compiler/InputConversion/MHLO/test/legalize_input_types.mlir
+++ b/iree/compiler/InputConversion/MHLO/test/legalize_input_types.mlir
@@ -149,9 +149,9 @@
 // CHECK-LABEL: func @linalg_non_structured_op
 // CHECK-SAME:    (%arg0: tensor<9xi32>) -> tensor<1x9xi32>
 func @linalg_non_structured_op(%arg0: tensor<9xi64>) -> tensor<1x9xi64> {
-  // CHECK:       %[[RES:.+]] = linalg.tensor_expand_shape %arg0 {{\[}}[0, 1]] : tensor<9xi32> into tensor<1x9xi32>
+  // CHECK:       %[[RES:.+]] = tensor.expand_shape %arg0 {{\[}}[0, 1]] : tensor<9xi32> into tensor<1x9xi32>
   // CHECK:       return %[[RES:.+]] : tensor<1x9xi32>
-  %0 = linalg.tensor_expand_shape %arg0 [[0, 1]] : tensor<9xi64> into tensor<1x9xi64>
+  %0 = tensor.expand_shape %arg0 [[0, 1]] : tensor<9xi64> into tensor<1x9xi64>
   return %0 : tensor<1x9xi64>
 }
 
diff --git a/iree/test/e2e/xla_ops/scatter.mlir b/iree/test/e2e/xla_ops/scatter.mlir
index 259a78e..2614c28 100644
--- a/iree/test/e2e/xla_ops/scatter.mlir
+++ b/iree/test/e2e/xla_ops/scatter.mlir
@@ -108,7 +108,7 @@
      %1 = arith.index_cast %0 : index to i32
      linalg.yield %1 : i32
       } -> tensor<1400xi32>
-  %indices_reshaped = linalg.tensor_expand_shape %indices [[0, 1]] :
+  %indices_reshaped = tensor.expand_shape %indices [[0, 1]] :
       tensor<1400xi32> into tensor<1400x1xi32>
   %result = "mhlo.scatter"(%original, %indices_reshaped, %update)({
     ^bb0(%arg3 : tensor<i32>, %arg4 : tensor<i32>):
@@ -139,7 +139,7 @@
         %1 = arith.index_cast %0 : index to i32
         linalg.yield %1 : i32
       } -> tensor<200xi32>
-  %indices_reshaped = linalg.tensor_expand_shape %indices [[0, 1]] :
+  %indices_reshaped = tensor.expand_shape %indices [[0, 1]] :
       tensor<200xi32> into tensor<200x1xi32>
   %result = "mhlo.scatter"(%original, %indices_reshaped, %update)({
     ^bb0(%arg3 : tensor<i32>, %arg4 : tensor<i32>):
diff --git a/iree/tools/BUILD b/iree/tools/BUILD
index 4be1159..4a29ef1 100644
--- a/iree/tools/BUILD
+++ b/iree/tools/BUILD
@@ -160,6 +160,7 @@
         "@llvm-project//mlir:Shape",
         "@llvm-project//mlir:StandardOps",
         "@llvm-project//mlir:StandardToSPIRV",
+        "@llvm-project//mlir:TensorInferTypeOpInterfaceImpl",
         "@llvm-project//mlir:Transforms",
         "@llvm-project//mlir:VectorOps",
     ],
diff --git a/iree/tools/init_mlir_dialects.h b/iree/tools/init_mlir_dialects.h
index 2ddcf82..fc8d911 100644
--- a/iree/tools/init_mlir_dialects.h
+++ b/iree/tools/init_mlir_dialects.h
@@ -25,6 +25,7 @@
 #include "mlir/Dialect/Shape/IR/Shape.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/Dialect/Tensor/IR/Tensor.h"
+#include "mlir/Dialect/Tensor/IR/TensorInferTypeOpInterfaceImpl.h"
 #include "mlir/Dialect/Tosa/IR/TosaOps.h"
 #include "mlir/Dialect/Vector/VectorOps.h"
 #include "mlir/IR/Dialect.h"
@@ -55,6 +56,7 @@
                   tosa::TosaDialect,
                   shape::ShapeDialect>();
   // clang-format on
+  tensor::registerInferTypeOpInterfaceExternalModels(registry);
 
 #ifdef IREE_HAVE_EMITC_DIALECT
   registry.insert<emitc::EmitCDialect>();
diff --git a/third_party/llvm-project b/third_party/llvm-project
index def8b95..c6a8bec 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit def8b952ebc00b0ad0fa4196ae27b11a385087ad
+Subproject commit c6a8bec4c578a7c12e4458b161fce7b1704804a2