diff --git a/bindings/python/iree/runtime/system_api_test.py b/bindings/python/iree/runtime/system_api_test.py
index 7976a12..fbe6006 100644
--- a/bindings/python/iree/runtime/system_api_test.py
+++ b/bindings/python/iree/runtime/system_api_test.py
@@ -64,10 +64,7 @@
     f = ctx.modules.arithmetic["simple_mul"]
     f_repr = repr(f)
     logging.info("f_repr: %s", f_repr)
-    self.assertRegex(
-        f_repr,
-        re.escape(
-            "(Buffer<float32[4]>, Buffer<float32[4]>) -> (Buffer<float32[4]>)"))
+    self.assertEqual(f_repr, "<VmFunction simple_mul(0rr_r), reflection = {}>")
 
   def test_duplicate_module(self):
     ctx = iree.runtime.SystemContext()
@@ -87,18 +84,19 @@
     results = f(arg0, arg1)
     np.testing.assert_allclose(results, [4., 10., 18., 28.])
 
-  def test_serialize_values(self):
-    ctx = iree.runtime.SystemContext()
-    self.assertTrue(ctx.is_dynamic)
-    ctx.add_vm_module(create_simple_mul_module())
-    self.assertEqual(ctx.modules.arithmetic.name, "arithmetic")
-    f = ctx.modules.arithmetic["simple_mul"]
-    arg0 = np.array([1., 2., 3., 4.], dtype=np.float32)
-    arg1 = np.array([4., 5., 6., 7.], dtype=np.float32)
-    results = f(arg0, arg1)
-    inputs, outputs = f.get_serialized_values()
-    self.assertEqual(inputs, ("4xf32=1 2 3 4", "4xf32=4 5 6 7"))
-    self.assertEqual(outputs, ("4xf32=4 10 18 28",))
+  # TODO: Re-implement tracing in a more sustainable fashion.
+  # def test_serialize_values(self):
+  #   ctx = iree.runtime.SystemContext()
+  #   self.assertTrue(ctx.is_dynamic)
+  #   ctx.add_vm_module(create_simple_mul_module())
+  #   self.assertEqual(ctx.modules.arithmetic.name, "arithmetic")
+  #   f = ctx.modules.arithmetic["simple_mul"]
+  #   arg0 = np.array([1., 2., 3., 4.], dtype=np.float32)
+  #   arg1 = np.array([4., 5., 6., 7.], dtype=np.float32)
+  #   results = f(arg0, arg1)
+  #   inputs, outputs = f.get_serialized_values()
+  #   self.assertEqual(inputs, ("4xf32=1 2 3 4", "4xf32=4 5 6 7"))
+  #   self.assertEqual(outputs, ("4xf32=4 10 18 28",))
 
   def test_load_vm_module(self):
     arithmetic = iree.runtime.load_vm_module(create_simple_mul_module())
diff --git a/bindings/python/tests/compiler_core_test.py b/bindings/python/tests/compiler_core_test.py
index 8c0df9a..5ccce84 100644
--- a/bindings/python/tests/compiler_core_test.py
+++ b/bindings/python/tests/compiler_core_test.py
@@ -158,7 +158,8 @@
         strip_source_map=True,
         strip_symbols=True,
         crash_reproducer_path="foobar.txt",
-        enable_benchmark=True,
+        # Re-enable when benchmarking pass is fixed: #6196
+        # enable_benchmark=True,
         target_backends=iree.compiler.DEFAULT_TESTING_BACKENDS)
 
   def testException(self):
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/BUILD b/integrations/tensorflow/iree_tf_compiler/TF/BUILD
index f241293..004ecd2 100644
--- a/integrations/tensorflow/iree_tf_compiler/TF/BUILD
+++ b/integrations/tensorflow/iree_tf_compiler/TF/BUILD
@@ -41,6 +41,7 @@
         "@iree//iree/compiler/Bindings/SIP/Utils",
         "@iree//iree/compiler/Dialect/Flow/IR",
         "@iree//iree/compiler/Dialect/Flow/Transforms",
+        "@iree//iree/compiler/Dialect/HAL/IR",
         "@iree//iree/compiler/Dialect/IREE/IR",
         "@iree//iree/compiler/Dialect/Shape/Transforms",
         "@llvm-project//llvm:Support",
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/Passes.cpp b/integrations/tensorflow/iree_tf_compiler/TF/Passes.cpp
index af3996c..f86d99d 100644
--- a/integrations/tensorflow/iree_tf_compiler/TF/Passes.cpp
+++ b/integrations/tensorflow/iree_tf_compiler/TF/Passes.cpp
@@ -112,11 +112,9 @@
   // - It assumes that tf_saved_model.bound_inputs have been eliminated
   // - It removes tf_saved_model.semantics from the module, which we can only
   //   do at the very end.
-  pm.addPass(createLowerExportedFunctionsPass());
-  // TODO: Remove the above and uncomment the below to enable IREE native ABI.
-  // pm.addPass(createSavedModelToIREEABIPass());
-  // // Inline the wrapper functions.
-  // pm.addPass(createInlinerPass());
+  pm.addPass(createSavedModelToIREEABIPass());
+  // Inline the wrapper functions.
+  pm.addPass(createInlinerPass());
 
   //----------------------------------------------------------------------------
   // Ensure that all Tensorflow has been legalized away
diff --git a/integrations/tensorflow/iree_tf_compiler/TF/SavedModelToIreeABI.cpp b/integrations/tensorflow/iree_tf_compiler/TF/SavedModelToIreeABI.cpp
index afeb805..98f81a5 100644
--- a/integrations/tensorflow/iree_tf_compiler/TF/SavedModelToIreeABI.cpp
+++ b/integrations/tensorflow/iree_tf_compiler/TF/SavedModelToIreeABI.cpp
@@ -18,6 +18,8 @@
 
 #include "iree/compiler/Dialect/Flow/IR/FlowDialect.h"
 #include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
+#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
+#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
 #include "iree/compiler/Dialect/IREE/IR/IREEDialect.h"
 #include "iree/compiler/Dialect/IREE/IR/IREEOps.h"
 #include "iree/compiler/Dialect/IREE/IR/IREETypes.h"
@@ -37,6 +39,7 @@
 #include "tensorflow/compiler/mlir/tensorflow/ir/tf_saved_model.h"
 
 namespace json = llvm::json;
+namespace IREE = mlir::iree_compiler::IREE;
 
 namespace mlir {
 namespace iree_integrations {
@@ -130,14 +133,16 @@
   }
 
   Type getIrType(Builder builder) {
-    auto variantType =
-        iree_compiler::IREE::VariantType::get(builder.getContext());
+    auto variantType = IREE::VariantType::get(builder.getContext());
     if (type == LevelType::Value) {
+      if (valueType.isa<TensorType>()) {
+        return IREE::HAL::BufferViewType::get(builder.getContext());
+      }
       return valueType;
     } else if (type == LevelType::List || type == LevelType::Tuple) {
-      return iree_compiler::IREE::ListType::get(variantType);
+      return IREE::ListType::get(variantType);
     } else if (type == LevelType::Dict) {
-      return iree_compiler::IREE::ListType::get(variantType);
+      return IREE::ListType::get(variantType);
     }
 
     llvm_unreachable("Unknown LevelType");
@@ -206,7 +211,12 @@
     if (type == LevelType::Value) {
       assert(valueIndex < callArgs.size() && "mismatched number of call args");
       assert(!callArgs[valueIndex] && "duplicate argument bindings");
-      callArgs[valueIndex] = thisValue;
+      auto value = thisValue;
+      if (value.getType().isa<IREE::HAL::BufferViewType>()) {
+        value = builder.createOrFold<IREE::HAL::TensorCastOp>(loc, valueType,
+                                                              thisValue);
+      }
+      callArgs[valueIndex] = value;
       return;
     }
 
@@ -241,23 +251,26 @@
     if (type == LevelType::Value) {
       assert(valueIndex < callReturns.size() &&
              "mismatched number of call returns");
-      return callReturns[valueIndex];
+      Value value = callReturns[valueIndex];
+      if (valueType.isa<TensorType>()) {
+        value = builder.createOrFold<IREE::HAL::TensorCastOp>(
+            loc, getIrType(builder), value);
+      }
+      return value;
     }
     // Recurse into sequence (index can be sparse on child ikey).
     if (type == LevelType::List || type == LevelType::Tuple) {
       Value listSizeValue =
           builder.create<ConstantOp>(loc, builder.getIndexType(),
                                      builder.getIndexAttr(getNeededListSize()));
-      Value listValue = builder.create<iree_compiler::IREE::ListCreateOp>(
+      Value listValue = builder.create<IREE::ListCreateOp>(
           loc, getIrType(builder), listSizeValue);
-      builder.create<iree_compiler::IREE::ListResizeOp>(loc, listValue,
-                                                        listSizeValue);
+      builder.create<IREE::ListResizeOp>(loc, listValue, listSizeValue);
       for (StructureLevel &child : children) {
         Value childValue = child.emitCreateReturns(loc, builder, callReturns);
         Value indexValue = builder.create<ConstantOp>(
             loc, builder.getIndexType(), builder.getIndexAttr(child.ikey));
-        builder.create<iree_compiler::IREE::ListSetOp>(loc, listValue,
-                                                       indexValue, childValue);
+        builder.create<IREE::ListSetOp>(loc, listValue, indexValue, childValue);
       }
       return listValue;
     }
@@ -267,17 +280,15 @@
       Value listSizeValue =
           builder.create<ConstantOp>(loc, builder.getIndexType(),
                                      builder.getIndexAttr(getNeededListSize()));
-      Value listValue = builder.create<iree_compiler::IREE::ListCreateOp>(
+      Value listValue = builder.create<IREE::ListCreateOp>(
           loc, getIrType(builder), listSizeValue);
-      builder.create<iree_compiler::IREE::ListResizeOp>(loc, listValue,
-                                                        listSizeValue);
+      builder.create<IREE::ListResizeOp>(loc, listValue, listSizeValue);
       for (auto it : llvm::enumerate(children)) {
         StructureLevel &child = it.value();
         Value childValue = child.emitCreateReturns(loc, builder, callReturns);
         Value indexValue = builder.create<ConstantOp>(
             loc, builder.getIndexType(), builder.getIndexAttr(it.index()));
-        builder.create<iree_compiler::IREE::ListSetOp>(loc, listValue,
-                                                       indexValue, childValue);
+        builder.create<IREE::ListSetOp>(loc, listValue, indexValue, childValue);
       }
       return listValue;
     }
@@ -290,10 +301,14 @@
                         int index) {
     Value indexValue = builder.create<ConstantOp>(loc, builder.getIndexType(),
                                                   builder.getIndexAttr(index));
-    Value itemValue = builder.create<iree_compiler::IREE::ListGetOp>(
-        loc, getIrType(builder), parentList, indexValue);
+    Value itemValue = builder.create<IREE::ListGetOp>(loc, getIrType(builder),
+                                                      parentList, indexValue);
     // TODO: Null check, etc. How does that work if returning a tensor? Need
     // to box somehow?
+    if (itemValue.getType().isa<IREE::HAL::BufferViewType>()) {
+      itemValue = builder.createOrFold<IREE::HAL::TensorCastOp>(loc, valueType,
+                                                                itemValue);
+    }
     return itemValue;
   }
 
@@ -548,8 +563,8 @@
     : public PassWrapper<SavedModelToIREEABIPass, OperationPass<ModuleOp>> {
  public:
   void getDependentDialects(DialectRegistry &registry) const override {
-    registry.insert<iree_compiler::IREE::Flow::FlowDialect,
-                    iree_compiler::IREEDialect,
+    registry.insert<IREE::Flow::FlowDialect, iree_compiler::IREEDialect,
+                    IREE::HAL::HALDialect,
                     mlir::tf_saved_model::TensorFlowSavedModelDialect>();
   }
 
diff --git a/integrations/tensorflow/iree_tf_compiler/iree-import-xla-main.cpp b/integrations/tensorflow/iree_tf_compiler/iree-import-xla-main.cpp
index dbff7f9..43abd66 100644
--- a/integrations/tensorflow/iree_tf_compiler/iree-import-xla-main.cpp
+++ b/integrations/tensorflow/iree_tf_compiler/iree-import-xla-main.cpp
@@ -260,9 +260,8 @@
 
   // Note that we emit the ABI last since any needed function-level
   // transformations (i.e. de-tupling, etc) should have been done.
-  // TODO: Uncomment this to enable IREE native bindings.
-  // pm.addNestedPass<FuncOp>(
-  //     iree_integrations::TF::createEmitDefaultIREEABIPass());
+  pm.addNestedPass<FuncOp>(
+      iree_integrations::MHLO::createEmitDefaultIREEABIPass());
 
   if (failed(pm.run(*module))) {
     llvm::errs()
diff --git a/integrations/tensorflow/iree_tf_compiler/iree-tf-opt-main.cpp b/integrations/tensorflow/iree_tf_compiler/iree-tf-opt-main.cpp
index 04e8d47..982b5ce 100644
--- a/integrations/tensorflow/iree_tf_compiler/iree-tf-opt-main.cpp
+++ b/integrations/tensorflow/iree_tf_compiler/iree-tf-opt-main.cpp
@@ -10,12 +10,20 @@
 // passes here. If you need something, add it, but add only what you need as
 // each addition will likely end up on the build critical path.
 
+#include "iree/compiler/Dialect/Flow/IR/FlowDialect.h"
+#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
+#include "iree/compiler/Dialect/IREE/IR/IREEDialect.h"
+#include "iree/compiler/Dialect/Modules/Strings/IR/Dialect.h"
+#include "iree/compiler/Dialect/Modules/TensorList/IR/TensorListOps.h"
 #include "iree/tools/init_xla_dialects.h"
 #include "iree_tf_compiler/MHLO/Passes.h"
 #include "iree_tf_compiler/TF/Passes.h"
+#include "iree_tf_compiler/dialect/tf_strings/ir/dialect.h"
+#include "iree_tf_compiler/dialect/tf_tensorlist/ir/tf_tensorlist_dialect.h"
 #include "llvm/Support/InitLLVM.h"
 #include "mlir/IR/Dialect.h"
 #include "mlir/Support/MlirOptMain.h"
+#include "mlir/Transforms/Passes.h"
 #include "tensorflow/compiler/mlir/tensorflow/dialect_registration.h"
 
 int main(int argc, char **argv) {
@@ -23,12 +31,25 @@
 
   mlir::DialectRegistry registry;
   mlir::registerXLADialects(registry);
+  registry.insert<mlir::iree_compiler::IREE::Flow::FlowDialect,
+                  mlir::iree_compiler::IREE::HAL::HALDialect,
+                  mlir::iree_compiler::IREEDialect,
+                  mlir::iree_compiler::IREE::Strings::StringsDialect,
+                  mlir::iree_compiler::IREE::TensorList::TensorListDialect>();
+  registry.insert<mlir::iree_integrations::tf_strings::TFStringsDialect>();
+  registry
+      .insert<mlir::iree_integrations::tf_tensorlist::TFTensorListDialect>();
 
   mlir::RegisterAllTensorFlowDialects(registry);
   mlir::iree_integrations::TF::registerAllDialects(registry);
   mlir::iree_integrations::TF::registerAllPasses();
   mlir::iree_integrations::MHLO::registerAllPasses();
 
+  // Select MLIR passes.
+  mlir::registerCanonicalizerPass();
+  mlir::registerCSEPass();
+  mlir::registerInlinerPass();
+
   if (failed(MlirOptMain(argc, argv, "IREE-TF modular optimizer driver\n",
                          registry,
                          /*preloadDialectsInContext=*/false))) {
diff --git a/iree/compiler/Dialect/Flow/Transforms/ExportBenchmarkFuncs.cpp b/iree/compiler/Dialect/Flow/Transforms/ExportBenchmarkFuncs.cpp
index e53e210..6a1eda6 100644
--- a/iree/compiler/Dialect/Flow/Transforms/ExportBenchmarkFuncs.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/ExportBenchmarkFuncs.cpp
@@ -49,9 +49,11 @@
                                                     OpBuilder& moduleBuilder) {
     std::string baseName = "_benchmark_input_";
     std::string name = baseName + std::to_string(uniqueId++);
-    auto variableOp = moduleBuilder.create<VariableOp>(
-        loc, name,
-        /*isMutable=*/false, inputType, moduleBuilder.getZeroAttr(inputType));
+    auto initialValue = moduleBuilder.getZeroAttr(inputType);
+    assert(initialValue && "failed to get zero attr for type");
+    auto variableOp = moduleBuilder.create<VariableOp>(loc, name,
+                                                       /*isMutable=*/false,
+                                                       inputType, initialValue);
     variableOp.setPrivate();
     variableOp->setAttr("noinline", UnitAttr::get(moduleBuilder.getContext()));
     return variableOp;
diff --git a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertStreamOps.cpp b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertStreamOps.cpp
index 6aabff7..d02ab8e 100644
--- a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertStreamOps.cpp
+++ b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/ConvertStreamOps.cpp
@@ -1018,8 +1018,7 @@
                   streamOp.getLoc(),
                   IREE::HAL::BufferType::get(rewriter.getContext()),
                   bufferValue),
-              rewriter.createOrFold<IREE::HAL::BufferViewByteLengthOp>(
-                  streamOp.getLoc(), bufferValue)};
+              schedulingState.lookupOrComputeSize(streamValue, rewriter)};
         } else {
           bufferRange = BufferRange{
               bufferValue,
diff --git a/iree/compiler/Dialect/HAL/Conversion/IREEToHAL/ConvertIREEToHAL.cpp b/iree/compiler/Dialect/HAL/Conversion/IREEToHAL/ConvertIREEToHAL.cpp
index 7bbf0ff..0320dfa 100644
--- a/iree/compiler/Dialect/HAL/Conversion/IREEToHAL/ConvertIREEToHAL.cpp
+++ b/iree/compiler/Dialect/HAL/Conversion/IREEToHAL/ConvertIREEToHAL.cpp
@@ -68,16 +68,60 @@
   }
 };
 
+template <typename T>
+class GenericConvertTypesConversion : public OpConversionPattern<T> {
+ public:
+  using OpConversionPattern<T>::OpConversionPattern;
+
+  LogicalResult matchAndRewrite(
+      T op, llvm::ArrayRef<Value> newOperands,
+      ConversionPatternRewriter &rewriter) const override {
+    SmallVector<Type> newTypes;
+    bool anyChanged = false;
+    for (auto oldNew : llvm::zip(op->getOperands(), newOperands)) {
+      auto oldValue = std::get<0>(oldNew);
+      auto newValue = std::get<1>(oldNew);
+      if (oldValue.getType() != newValue.getType()) {
+        anyChanged = true;
+        break;
+      }
+    }
+    for (auto oldType : op.getOperation()->getResultTypes()) {
+      auto newType = this->getTypeConverter()->convertType(oldType);
+      if (oldType != newType) anyChanged = true;
+      newTypes.push_back(newType);
+    }
+    if (!anyChanged) return failure();
+    rewriter.replaceOpWithNewOp<T>(op, newTypes, newOperands, op->getAttrs());
+    return success();
+  }
+};
+
 }  // namespace
 
-// Appends all patterns for lowering IREE ops to HAL buffer ops.
-void populateIREEToHALPatterns(MLIRContext *context,
+void populateIREEToHALPatterns(MLIRContext *context, ConversionTarget &target,
+                               TypeConverter &typeConverter,
                                OwningRewritePatternList &patterns) {
-  patterns.insert<DynamicShapeConstantOpConversion>(context);
-}
-
-void setupIREEToHALLegality(MLIRContext *context, ConversionTarget &target) {
   target.addIllegalOp<IREE::DynamicShapeConstantOp>();
+  patterns.insert<DynamicShapeConstantOpConversion>(context);
+
+  typeConverter.addConversion([&](IREE::ListType type) {
+    auto elementType = typeConverter.convertType(type.getElementType());
+    return IREE::ListType::get(elementType);
+  });
+
+  target.addDynamicallyLegalOp<IREE::ListCreateOp>([&](IREE::ListCreateOp op) {
+    return typeConverter.isLegal(op.getType());
+  });
+  target.addDynamicallyLegalOp<IREE::ListGetOp>(
+      [&](IREE::ListGetOp op) { return typeConverter.isLegal(op.getType()); });
+  target.addDynamicallyLegalOp<IREE::ListSetOp>([&](IREE::ListSetOp op) {
+    return typeConverter.isLegal(op.value().getType());
+  });
+  patterns.insert<GenericConvertTypesConversion<IREE::ListCreateOp>,
+                  GenericConvertTypesConversion<IREE::ListGetOp>,
+                  GenericConvertTypesConversion<IREE::ListSetOp>>(typeConverter,
+                                                                  context);
 }
 
 }  // namespace iree_compiler
diff --git a/iree/compiler/Dialect/HAL/Conversion/IREEToHAL/ConvertIREEToHAL.h b/iree/compiler/Dialect/HAL/Conversion/IREEToHAL/ConvertIREEToHAL.h
index c643ab7..128110d 100644
--- a/iree/compiler/Dialect/HAL/Conversion/IREEToHAL/ConvertIREEToHAL.h
+++ b/iree/compiler/Dialect/HAL/Conversion/IREEToHAL/ConvertIREEToHAL.h
@@ -17,14 +17,12 @@
 // TODO(gcmn): Use conversion interfaces. Requires breaking circular dependency
 // between HAL and IREE dialects.
 
-// Appends all patterns for lowering IREE ops to HAL buffer ops.
-void populateIREEToHALPatterns(MLIRContext *context,
+// Appends all patterns for lowering IREE ops to HAL buffer ops and sets their
+// legality.
+void populateIREEToHALPatterns(MLIRContext *context, ConversionTarget &target,
+                               TypeConverter &typeConverter,
                                OwningRewritePatternList &patterns);
 
-// Setup the |conversionTarget| op legality to ensure helpful error messages for
-// IREE ops we know should always be converted.
-void setupIREEToHALLegality(MLIRContext *context, ConversionTarget &target);
-
 }  // namespace iree_compiler
 }  // namespace mlir
 
diff --git a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStandardToHAL.cpp b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStandardToHAL.cpp
index f0510e2..34fa2f9 100644
--- a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStandardToHAL.cpp
+++ b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStandardToHAL.cpp
@@ -24,23 +24,42 @@
  public:
   using OpConversionPattern::OpConversionPattern;
   LogicalResult matchAndRewrite(
-      IREE::HAL::TensorCastOp op, llvm::ArrayRef<Value> operands,
+      IREE::HAL::TensorCastOp op, llvm::ArrayRef<Value> rawOperands,
       ConversionPatternRewriter &rewriter) const override {
+    IREE::HAL::TensorCastOpAdaptor newOperands(
+        rawOperands, op.getOperation()->getAttrDictionary());
     Value newValue = {};
     auto targetType = op.target().getType();
     if (targetType.isa<TensorType>()) {
       // HAL type -> tensor<...>
-      newValue = operands.front();
+      newValue = newOperands.source();
     } else if (targetType.isa<IREE::HAL::BufferType>()) {
       // tensor<...> -> !hal.buffer
       auto adaptor = IREE::HAL::TensorRewriteAdaptor::get(
-          op.getLoc(), op.source(), operands.front(), rewriter);
+          op.getLoc(), op.source(), newOperands.source(), rewriter);
       newValue = adaptor.getBuffer();
     } else if (targetType.isa<IREE::HAL::BufferViewType>()) {
       // tensor<...> -> !hal.buffer_view
       auto adaptor = IREE::HAL::TensorRewriteAdaptor::get(
-          op.getLoc(), op.source(), operands.front(), rewriter);
-      newValue = adaptor.getBufferView();
+          op.getLoc(), op.source(), newOperands.source(), rewriter);
+
+      // Note that the buffer view cannot just be returned here: it's backing
+      // buffer will be correct, but the cast may be doing a metadata change,
+      // which must be reflected in the returned buffer view. For now, we
+      // just create a new view unconditionally when converting from a tensor
+      // since that is conservative. But this can be optimized with additional
+      // heuristics regarding when it is safe to alias the original.
+      Value originalValue = op.source();
+      if (auto sourceType =
+              originalValue.getType().dyn_cast<RankedTensorType>()) {
+        auto shapeDims = getShapeDims(rewriter, op.getLoc(), sourceType,
+                                      newOperands.source_dims());
+        newValue = rewriter.create<IREE::HAL::BufferViewCreateOp>(
+            op.getLoc(), adaptor.getBuffer(), adaptor.getElementType(),
+            shapeDims);
+      } else {
+        newValue = adaptor.getBufferView();
+      }
     }
     if (!newValue) {
       return rewriter.notifyMatchFailure(op, "bad source/target type pair");
@@ -48,6 +67,22 @@
     rewriter.replaceOp(op, {newValue});
     return success();
   }
+
+  SmallVector<Value> getShapeDims(OpBuilder &builder, Location loc,
+                                  RankedTensorType sourceType,
+                                  ValueRange sourceDims) const {
+    SmallVector<Value> shapeDims(sourceType.getRank());
+    int sourceDimsIndex = 0;
+    for (int i = 0, e = shapeDims.size(); i < e; ++i) {
+      if (sourceType.isDynamicDim(i)) {
+        shapeDims[i] = sourceDims[sourceDimsIndex++];
+      } else {
+        shapeDims[i] =
+            builder.create<ConstantIndexOp>(loc, sourceType.getDimSize(i));
+      }
+    }
+    return shapeDims;
+  }
 };
 
 }  // namespace
diff --git a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/test/structural_ops.mlir b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/test/structural_ops.mlir
index 5464507..01665d1 100644
--- a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/test/structural_ops.mlir
+++ b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/test/structural_ops.mlir
@@ -19,3 +19,29 @@
   // CHECK: return %[[RES]] : !hal.buffer
   return %0 : tensor<i32>
 }
+
+// -----
+// CHECK-LABEL: func @tensor_cast_does_not_alias_metadata_update
+func @tensor_cast_does_not_alias_metadata_update(%arg0: !hal.buffer_view) -> !hal.buffer_view {
+    %0 = hal.tensor.cast %arg0 : !hal.buffer_view -> tensor<3x2x2x2xf32>
+    %1 = flow.ex.stream.fragment(%0) : (tensor<3x2x2x2xf32>) -> (tensor<3x2x1x4x1xf32>) =
+        (%arg1: tensor<3x2x2x2xf32>) -> tensor<3x2x1x4x1xf32> {
+      %3 = flow.tensor.reshape %arg1 : tensor<3x2x2x2xf32> -> tensor<3x2x1x4x1xf32>
+      flow.return %3 : tensor<3x2x1x4x1xf32>
+    }
+
+    // Just anchor on the end of the function that creates a new buffer view.
+    // CHECK: hal.ex.submit_and_wait
+    // CHECK: %[[C3:.*]] = constant 3 : index
+    // CHECK: %[[C2:.*]] = constant 2 : index
+    // CHECK: %[[C1_1:.*]] = constant 1 : index
+    // CHECK: %[[C4:.*]] = constant 4 : index
+    // CHECK: %[[C1_2:.*]] = constant 1 : index
+    // CHECK: %[[ET:.*]] = constant {{.*}} : i32
+    // CHECK: %[[VIEW:.*]] = hal.buffer_view.create
+    //   CHECK-SAME: element_type = %[[ET]],
+    //   CHECK-SAME: shape = [%[[C3]], %[[C2]], %[[C1_1]], %[[C4]], %[[C1_2]]]
+    // CHECK: return %[[VIEW]]
+    %2 = hal.tensor.cast %1 : tensor<3x2x1x4x1xf32> -> !hal.buffer_view
+    return %2 : !hal.buffer_view
+}
diff --git a/iree/compiler/Dialect/HAL/Conversion/TypeConverter.cpp b/iree/compiler/Dialect/HAL/Conversion/TypeConverter.cpp
index 9fe322e..ce5d4f2 100644
--- a/iree/compiler/Dialect/HAL/Conversion/TypeConverter.cpp
+++ b/iree/compiler/Dialect/HAL/Conversion/TypeConverter.cpp
@@ -29,12 +29,11 @@
   });
 
   // Tensors become buffers by default.
-  // TODO(benvanik): make them buffer views instead? then they carry shape but
-  // are memory type erased which is not good.
+  // Shapes and types are carried independently or folded away entirely - all
+  // we need at the HAL level is a blob of bytes.
   addConversion([](TensorType type) -> Optional<Type> {
     // HAL only should be concerned with numeric values.
     if (HALTypeConverter::shouldConvertToBuffer(type)) {
-      // TODO(benvanik): composite-type conversion (buffer + dynamic dims).
       return IREE::HAL::BufferType::get(type.getContext());
     }
     return llvm::None;
@@ -58,8 +57,22 @@
                               IREE::HAL::BufferViewType type, ValueRange inputs,
                               Location loc) -> Value {
     assert(inputs.size() == 1);
-    assert(inputs[0].getType().isa<TensorType>());
-    return builder.create<IREE::HAL::TensorCastOp>(loc, type, inputs[0]);
+    auto inputValue = inputs[0];
+    auto inputType = inputValue.getType();
+    if (inputType.isa<TensorType>()) {
+      return builder.create<IREE::HAL::TensorCastOp>(loc, type, inputValue);
+    } else if (inputType.isa<IREE::HAL::BufferType>()) {
+      // Look for the buffer view this buffer came from, if any.
+      // If we don't have the origin buffer view then we can't know the shape
+      // and can't materialize one here - it's too late.
+      if (auto bvbOp = dyn_cast_or_null<IREE::HAL::BufferViewBufferOp>(
+              inputValue.getDefiningOp())) {
+        return bvbOp.buffer_view();
+      }
+      return nullptr;
+    } else {
+      return nullptr;
+    }
   });
 
   // Recursively handle pointer target types (we want to convert
diff --git a/iree/compiler/Dialect/HAL/Transforms/ConvertToHAL.cpp b/iree/compiler/Dialect/HAL/Transforms/ConvertToHAL.cpp
index 513536c..959ac6e 100644
--- a/iree/compiler/Dialect/HAL/Transforms/ConvertToHAL.cpp
+++ b/iree/compiler/Dialect/HAL/Transforms/ConvertToHAL.cpp
@@ -65,8 +65,8 @@
 
     OwningRewritePatternList patterns(&getContext());
 
-    setupIREEToHALLegality(context, conversionTarget);
-    populateIREEToHALPatterns(context, patterns);
+    populateIREEToHALPatterns(context, conversionTarget, typeConverter,
+                              patterns);
 
     setupCompilerHintsLegality(context, conversionTarget, typeConverter);
     populatePreserveCompilerHintsPatterns(context, patterns);
diff --git a/iree/compiler/Dialect/Modules/Strings/IR/Dialect.cc b/iree/compiler/Dialect/Modules/Strings/IR/Dialect.cc
index da807a7..8fe8134 100644
--- a/iree/compiler/Dialect/Modules/Strings/IR/Dialect.cc
+++ b/iree/compiler/Dialect/Modules/Strings/IR/Dialect.cc
@@ -18,6 +18,7 @@
 #include "mlir/IR/SymbolTable.h"
 #include "mlir/Parser.h"
 #include "mlir/Transforms/DialectConversion.h"
+#include "mlir/Transforms/InliningUtils.h"
 
 namespace mlir {
 namespace iree_compiler {
@@ -26,6 +27,28 @@
 
 namespace {
 
+struct StringsInlinerInterface : public DialectInlinerInterface {
+  using DialectInlinerInterface::DialectInlinerInterface;
+
+  bool isLegalToInline(Operation *call, Operation *callable,
+                       bool wouldBeCloned) const final {
+    // Sure!
+    return true;
+  }
+
+  bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
+                       BlockAndValueMapping &valueMapping) const final {
+    // Sure!
+    return true;
+  }
+
+  bool isLegalToInline(Operation *op, Region *dest, bool wouldBeCloned,
+                       BlockAndValueMapping &valueMapping) const final {
+    // Sure!
+    return true;
+  }
+};
+
 class StringsToVMConversionInterface : public VMConversionDialectInterface {
  public:
   using VMConversionDialectInterface::VMConversionDialectInterface;
@@ -63,6 +86,7 @@
     : Dialect(getDialectNamespace(), context, TypeID::get<StringsDialect>()) {
   addInterfaces<StringsToVMConversionInterface>();
   addInterfaces<StringsToHALConversionInterface>();
+  addInterfaces<StringsInlinerInterface>();
 
   addTypes<StringType, StringTensorType>();
 
diff --git a/iree/compiler/Dialect/Modules/TensorList/IR/TensorListDialect.cpp b/iree/compiler/Dialect/Modules/TensorList/IR/TensorListDialect.cpp
index c07a4cf..9975127 100644
--- a/iree/compiler/Dialect/Modules/TensorList/IR/TensorListDialect.cpp
+++ b/iree/compiler/Dialect/Modules/TensorList/IR/TensorListDialect.cpp
@@ -17,6 +17,7 @@
 #include "mlir/IR/SymbolTable.h"
 #include "mlir/Parser.h"
 #include "mlir/Transforms/DialectConversion.h"
+#include "mlir/Transforms/InliningUtils.h"
 
 namespace mlir {
 namespace iree_compiler {
@@ -25,6 +26,28 @@
 
 namespace {
 
+struct TensorListInlinerInterface : public DialectInlinerInterface {
+  using DialectInlinerInterface::DialectInlinerInterface;
+
+  bool isLegalToInline(Operation *call, Operation *callable,
+                       bool wouldBeCloned) const final {
+    // Sure!
+    return true;
+  }
+
+  bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
+                       BlockAndValueMapping &valueMapping) const final {
+    // Sure!
+    return true;
+  }
+
+  bool isLegalToInline(Operation *op, Region *dest, bool wouldBeCloned,
+                       BlockAndValueMapping &valueMapping) const final {
+    // Sure!
+    return true;
+  }
+};
+
 class TensorListToHALConversionInterface
     : public HALConversionDialectInterface {
  public:
@@ -64,6 +87,7 @@
               TypeID::get<TensorListDialect>()) {
   addInterfaces<TensorListToHALConversionInterface>();
   addInterfaces<TensorListToVMConversionInterface>();
+  addInterfaces<TensorListInlinerInterface>();
 
   addTypes<TensorListType>();
 
diff --git a/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp b/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp
index deed07d..df29a52 100644
--- a/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp
+++ b/iree/compiler/Dialect/VM/IR/VMOpFolders.cpp
@@ -959,12 +959,12 @@
     // x / 1 = x
     return op.lhs();
   }
-  return constFoldBinaryOp<FloatAttr>(operands,
-                                      [](const APFloat &a, const APFloat &b) {
-                                        APFloat c = a;
-                                        c.divide(b, APFloat::rmTowardZero);
-                                        return c;
-                                      });
+  return constFoldBinaryOp<FloatAttr>(
+      operands, [](const APFloat &a, const APFloat &b) {
+        APFloat c = a;
+        c.divide(b, APFloat::rmNearestTiesToAway);
+        return c;
+      });
 }
 
 OpFoldResult DivF32Op::fold(ArrayRef<Attribute> operands) {
@@ -1012,7 +1012,7 @@
   return constFoldTernaryOp<FloatAttr>(
       operands, [](const APFloat &a, const APFloat &b, const APFloat &c) {
         APFloat d = a;
-        d.fusedMultiplyAdd(b, c, APFloat::rmTowardZero);
+        d.fusedMultiplyAdd(b, c, APFloat::rmNearestTiesToAway);
         return d;
       });
 }
@@ -1458,7 +1458,7 @@
   return constFoldCastOp<IntegerAttr, FloatAttr>(
       Float32Type::get(getContext()), operands, [&](const APInt &a) {
         APFloat b = APFloat(0.0f);
-        b.convertFromAPInt(a, /*IsSigned=*/true, APFloat::rmTowardZero);
+        b.convertFromAPInt(a, /*IsSigned=*/true, APFloat::rmNearestTiesToAway);
         return b;
       });
 }
@@ -1467,7 +1467,7 @@
   return constFoldCastOp<IntegerAttr, FloatAttr>(
       Float32Type::get(getContext()), operands, [&](const APInt &a) {
         APFloat b = APFloat(0.0f);
-        b.convertFromAPInt(a, /*IsSigned=*/false, APFloat::rmTowardZero);
+        b.convertFromAPInt(a, /*IsSigned=*/false, APFloat::rmNearestTiesToAway);
         return b;
       });
 }
@@ -1477,7 +1477,7 @@
       IntegerType::get(getContext(), 32), operands, [&](const APFloat &a) {
         bool isExact = false;
         llvm::APSInt b;
-        a.convertToInteger(b, APFloat::rmTowardZero, &isExact);
+        a.convertToInteger(b, APFloat::rmNearestTiesToAway, &isExact);
         return b;
       });
 }
@@ -1487,7 +1487,7 @@
       IntegerType::get(getContext(), 32), operands, [&](const APFloat &a) {
         bool isExact = false;
         llvm::APSInt b;
-        a.convertToInteger(b, APFloat::rmTowardZero, &isExact);
+        a.convertToInteger(b, APFloat::rmNearestTiesToAway, &isExact);
         b.setIsUnsigned(true);
         return b;
       });
diff --git a/iree/compiler/Translation/IREEVM.cpp b/iree/compiler/Translation/IREEVM.cpp
index ef86d73..e25b1a1 100644
--- a/iree/compiler/Translation/IREEVM.cpp
+++ b/iree/compiler/Translation/IREEVM.cpp
@@ -34,10 +34,10 @@
 // whole end-to-end with options for bindings/targets/etc.
 struct BindingOptions {
   // Whether to include runtime support functions for the IREE native ABI.
-  bool native = false;
+  bool native = true;
   // Whether to include runtime support functions and metadata required for
   // SIP-compatible bindings (like bindings/python/iree).
-  bool sip = true;
+  bool sip = false;
   // Whether to include runtime support functions required for the IREE TFLite
   // API compatibility bindings.
   bool tflite = false;
@@ -51,12 +51,12 @@
       "iree-native-bindings-support",
       llvm::cl::desc(
           "Include runtime support for native IREE ABI-compatible bindings"),
-      llvm::cl::init(false), llvm::cl::cat(bindingOptionsCategory)};
+      llvm::cl::init(true), llvm::cl::cat(bindingOptionsCategory)};
 
   static llvm::cl::opt<bool> *bindingsSIPFlag = new llvm::cl::opt<bool>{
       "iree-sip-bindings-support",
       llvm::cl::desc("Include runtime support for SIP-compatible bindings"),
-      llvm::cl::init(true), llvm::cl::cat(bindingOptionsCategory)};
+      llvm::cl::init(false), llvm::cl::cat(bindingOptionsCategory)};
 
   static llvm::cl::opt<bool> *bindingsTFLiteFlag = new llvm::cl::opt<bool>{
       "iree-tflite-bindings-support",
diff --git a/iree/samples/custom_modules/custom_modules_test.cc b/iree/samples/custom_modules/custom_modules_test.cc
index fbaec7d..9667784 100644
--- a/iree/samples/custom_modules/custom_modules_test.cc
+++ b/iree/samples/custom_modules/custom_modules_test.cc
@@ -129,22 +129,26 @@
 
 TEST_F(CustomModulesTest, PrintTensor) {
   // Allocate the buffer we'll be printing.
+  static iree_hal_dim_t kShape[] = {2, 4};
   static float kBufferContents[2 * 4] = {0.0f, 1.0f, 2.0f, 3.0f,
                                          4.0f, 5.0f, 6.0f, 7.0f};
-  iree_hal_buffer_t* buffer = nullptr;
-  IREE_ASSERT_OK(iree_hal_allocator_wrap_buffer(
-      hal_allocator_, IREE_HAL_MEMORY_TYPE_HOST_LOCAL,
+  iree_hal_buffer_view_t* buffer_view = nullptr;
+  IREE_ASSERT_OK(iree_hal_buffer_view_wrap_or_clone_heap_buffer(
+      hal_allocator_, kShape, IREE_ARRAYSIZE(kShape),
+      IREE_HAL_ELEMENT_TYPE_FLOAT_32,
+      IREE_HAL_MEMORY_TYPE_HOST_LOCAL | IREE_HAL_MEMORY_TYPE_DEVICE_VISIBLE,
       IREE_HAL_MEMORY_ACCESS_ALL, IREE_HAL_BUFFER_USAGE_ALL,
-      iree_byte_span_t{reinterpret_cast<uint8_t*>(kBufferContents),
-                       sizeof(kBufferContents)},
-      iree_allocator_null(), &buffer));
+      iree_make_byte_span((void*)kBufferContents, sizeof(kBufferContents)),
+      iree_allocator_null(), &buffer_view));
 
   // Pass in the tensor as an expanded HAL buffer.
   iree::vm::ref<iree_vm_list_t> inputs;
   IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1,
                                      iree_allocator_system(), &inputs));
-  iree_vm_ref_t input_buffer_ref = iree_hal_buffer_move_ref(buffer);
-  IREE_ASSERT_OK(iree_vm_list_push_ref_move(inputs.get(), &input_buffer_ref));
+  iree_vm_ref_t input_buffer_view_ref =
+      iree_hal_buffer_view_move_ref(buffer_view);
+  IREE_ASSERT_OK(
+      iree_vm_list_push_ref_move(inputs.get(), &input_buffer_view_ref));
 
   // Prepare outputs list to accept the results from the invocation.
   iree::vm::ref<iree_vm_list_t> outputs;
@@ -169,22 +173,26 @@
 
 TEST_F(CustomModulesTest, RoundTripTensor) {
   // Allocate the buffer we'll be printing/parsing.
+  static iree_hal_dim_t kShape[] = {2, 4};
   static float kBufferContents[2 * 4] = {0.0f, 1.0f, 2.0f, 3.0f,
                                          4.0f, 5.0f, 6.0f, 7.0f};
-  iree_hal_buffer_t* buffer = nullptr;
-  IREE_ASSERT_OK(iree_hal_allocator_wrap_buffer(
-      hal_allocator_, IREE_HAL_MEMORY_TYPE_HOST_LOCAL,
+  iree_hal_buffer_view_t* buffer_view = nullptr;
+  IREE_ASSERT_OK(iree_hal_buffer_view_wrap_or_clone_heap_buffer(
+      hal_allocator_, kShape, IREE_ARRAYSIZE(kShape),
+      IREE_HAL_ELEMENT_TYPE_FLOAT_32,
+      IREE_HAL_MEMORY_TYPE_HOST_LOCAL | IREE_HAL_MEMORY_TYPE_DEVICE_VISIBLE,
       IREE_HAL_MEMORY_ACCESS_ALL, IREE_HAL_BUFFER_USAGE_ALL,
-      iree_byte_span_t{reinterpret_cast<uint8_t*>(kBufferContents),
-                       sizeof(kBufferContents)},
-      iree_allocator_null(), &buffer));
+      iree_make_byte_span((void*)kBufferContents, sizeof(kBufferContents)),
+      iree_allocator_null(), &buffer_view));
 
   // Pass in the tensor as an expanded HAL buffer.
   iree::vm::ref<iree_vm_list_t> inputs;
   IREE_ASSERT_OK(iree_vm_list_create(/*element_type=*/nullptr, 1,
                                      iree_allocator_system(), &inputs));
-  iree_vm_ref_t input_buffer_ref = iree_hal_buffer_move_ref(buffer);
-  IREE_ASSERT_OK(iree_vm_list_push_ref_move(inputs.get(), &input_buffer_ref));
+  iree_vm_ref_t input_buffer_view_ref =
+      iree_hal_buffer_view_move_ref(buffer_view);
+  IREE_ASSERT_OK(
+      iree_vm_list_push_ref_move(inputs.get(), &input_buffer_view_ref));
 
   // Prepare outputs list to accept the results from the invocation.
   iree::vm::ref<iree_vm_list_t> outputs;
diff --git a/iree/test/e2e/regression/BUILD b/iree/test/e2e/regression/BUILD
index 2a78265..d77a6b2 100644
--- a/iree/test/e2e/regression/BUILD
+++ b/iree/test/e2e/regression/BUILD
@@ -28,7 +28,6 @@
         [
             "dynamic_abs.mlir",
             "dynamic_add.mlir",
-            "dynamic_compare_and_select.mlir",
             "dynamic_dot.mlir",
             "dynamic_torch_index_select_high_rank.mlir",
             "dynamic_torch_index_select_negative.mlir",
@@ -43,6 +42,7 @@
             ["*.mlir"],
         # TODO(#5897): enable these for codegen linalg on tensors/etc.
         exclude = [
+            "dynamic_compare_and_select.mlir",
             "dynamic_dot_general.mlir",
             "dynamic_linalg_matmul_on_tensors.mlir",
             "dynamic_linalg_matmul_on_tensors_fuse_0.mlir",
diff --git a/iree/test/e2e/regression/CMakeLists.txt b/iree/test/e2e/regression/CMakeLists.txt
index 4515c48..f01d088 100644
--- a/iree/test/e2e/regression/CMakeLists.txt
+++ b/iree/test/e2e/regression/CMakeLists.txt
@@ -16,7 +16,6 @@
   SRCS
     "dynamic_abs.mlir"
     "dynamic_add.mlir"
-    "dynamic_compare_and_select.mlir"
     "dynamic_dot.mlir"
     "dynamic_torch_index_select_high_rank.mlir"
     "dynamic_torch_index_select_negative.mlir"
diff --git a/iree/test/e2e/xla_ops/convert.mlir b/iree/test/e2e/xla_ops/convert.mlir
index 6a8606a..cee8b64 100644
--- a/iree/test/e2e/xla_ops/convert.mlir
+++ b/iree/test/e2e/xla_ops/convert.mlir
@@ -47,11 +47,15 @@
   return
 }
 
+// TODO(#6160): XLA does not specify the rounding behavior, meaning that we
+// can't test something like -10.5 as that could be -11 (roundf) or -10 (rint
+// with round-to-even mode).
+//
 // For casting rules, see
 // https://www.tensorflow.org/xla/operation_semantics#convertelementtype
-func @float_to_int() {
-  %input = iree.unfoldable_constant dense<[-10.5, -4.4, 4.4, 10.5]> : tensor<4xf32>
-  %res = "mhlo.convert"(%input) : (tensor<4xf32>) -> tensor<4xi32>
-  check.expect_eq_const(%res, dense<[-10, -4, 4, 10]> : tensor<4xi32>) : tensor<4xi32>
-  return
-}
+// func @float_to_int() {
+//   %input = iree.unfoldable_constant dense<[-10.5, -4.4, 4.4, 10.5]> : tensor<4xf32>
+//   %res = "mhlo.convert"(%input) : (tensor<4xf32>) -> tensor<4xi32>
+//   check.expect_eq_const(%res, dense<[-10, -4, 4, 10]> : tensor<4xi32>) : tensor<4xi32>
+//   return
+// }
diff --git a/iree/vm/ops.h b/iree/vm/ops.h
index 2cd3d3b..143ab41 100644
--- a/iree/vm/ops.h
+++ b/iree/vm/ops.h
@@ -99,10 +99,10 @@
   return (float)(uint32_t)operand;
 }
 static inline int32_t vm_cast_f32si32(float operand) {
-  return (int32_t)roundf(operand);
+  return (int32_t)lroundf(operand);
 }
 static inline int32_t vm_cast_f32ui32(float operand) {
-  return (uint32_t)roundf(operand);
+  return (uint32_t)lroundf(operand);
 }
 
 static inline float vm_atan_f32(float operand) { return atanf(operand); }