diff --git a/.github/workflows/lint.yml b/.github/workflows/lint.yml
index ad9635a..d5d5bb0 100644
--- a/.github/workflows/lint.yml
+++ b/.github/workflows/lint.yml
@@ -59,9 +59,10 @@
         run: |
           git diff -U0 "${BASE_REF?}" | python3 third_party/format_diff/format_diff.py yapf -i
           git diff --exit-code
-      - name: Instructions for fixing these linting errors
+      - name: Instructions for fixing the above linting errors
+        if: ${{ failure() }}
         run: |
-          printf "If the lint above failed it can be fixed by running\n"
+          printf "You can fix the lint errors above by running\n"
           printf "  git diff -U0 "${BASE_REF?}" | python3 third_party/format_diff/format_diff.py yapf -i\n"
 
   clang-format:
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 12a72fe..b2c30b7 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -96,8 +96,7 @@
   set( IREE_HAL_DRIVERS_TO_BUILD ${IREE_ALL_HAL_DRIVERS} )
   # For cross compilation towords Android, we don't want LLVM JIT HAL driver.
   if(ANDROID)
-    # TODO(ataei): Enable dylib/dylib-llvm-aot for android.
-    list(REMOVE_ITEM IREE_HAL_DRIVERS_TO_BUILD LLVM DyLib)
+    list(REMOVE_ITEM IREE_HAL_DRIVERS_TO_BUILD LLVM)
   endif()
 
   # For Apple platforms we need to use Metal instead of Vulkan.
diff --git a/SUBMODULE_VERSIONS b/SUBMODULE_VERSIONS
index 2aac894..e9e0de9 100644
--- a/SUBMODULE_VERSIONS
+++ b/SUBMODULE_VERSIONS
@@ -5,7 +5,7 @@
 a5d9d0f7d368054fd1691aedf1db4116efcc233e third_party/flatbuffers
 4fb0ff7069bd88ee85902f4d0bb62794e5f6d021 third_party/flatcc
 f2fb48c3b3d79a75a88a99fba6576b25d42ec528 third_party/googletest
-51ff04567b2f8d06b2062bd3ed72eab2e93e4466 third_party/llvm-project
+f402e682d0ef5598eeffc9a21a691b03e602ff58 third_party/llvm-project
 17b12a4481daa150e2d1ea3ada086b551b856707 third_party/marl
 d2cdb70e038370b5e28f353fe98ccd70af1cbc25 third_party/mlir-emitc
 d8c7ee00a687ac369e62e2032514a93a9b413502 third_party/pybind11
@@ -14,7 +14,7 @@
 685f86471e9d26b3eb7676695a2e2cefb4551ae9 third_party/spirv_cross
 f8bf11a0253a32375c32cad92c841237b96696c0 third_party/spirv_headers
 57eb48aed36160c4876bc8310d9ca84d42ee9e2a third_party/swiftshader
-74747f0c6b017df8fe87d1912f8eb4e2e287fbda third_party/tensorflow
+2d6bdab3adb0b8949763d5c63426338f938c9efe third_party/tensorflow
 a9a09ab0940408898fccfdcfe2bb8dc19b50f13c third_party/tracy
 9bd3f561bcee3f01d22912de10bb07ce4e23d378 third_party/vulkan_headers
 909f36b714c9239ee0b112a321220213a474ba53 third_party/vulkan_memory_allocator
diff --git a/build_tools/bazel/iree.bazelrc b/build_tools/bazel/iree.bazelrc
index 46922a4..0f246e7c 100644
--- a/build_tools/bazel/iree.bazelrc
+++ b/build_tools/bazel/iree.bazelrc
@@ -298,10 +298,10 @@
 # Another TensorFlow flag from their config script.
 build:windows --define with_default_optimizations=true
 
-# TensorFlow builds depend on these flags so we will as well
+# TensorFlow builds depend on this flag, but it doesn't appear to work with
+# gmock in some of our unit tests, so only enable it for TensorFlow files.
 # MSVC (Windows): Standards-conformant preprocessor mode
 # See https://docs.microsoft.com/en-us/cpp/preprocessor/preprocessor-experimental-overview
-build:windows --copt=/experimental:preprocessor
-build:windows --host_copt=/experimental:preprocessor
+build:windows --per_file_copt=tensorflow@/experimental:preprocessor
 
 ###############################################################################
diff --git a/build_tools/bazel/run_binary_test.bzl b/build_tools/bazel/run_binary_test.bzl
index 79536f3..aa165ea 100644
--- a/build_tools/bazel/run_binary_test.bzl
+++ b/build_tools/bazel/run_binary_test.bzl
@@ -32,12 +32,10 @@
 """
 
 def _run_binary_test_impl(ctx):
-    ctx.actions.run_shell(
-        inputs = [ctx.file.test_binary],
-        outputs = [ctx.outputs.executable],
-        command = "cp $1 $2",
-        arguments = [ctx.file.test_binary.path, ctx.outputs.executable.path],
-        mnemonic = "CopyExecutable",
+    ctx.actions.symlink(
+        target_file = ctx.executable.test_binary,
+        output = ctx.outputs.executable,
+        is_executable = True,
     )
 
     data_runfiles = ctx.runfiles(files = ctx.files.data)
@@ -54,7 +52,8 @@
     attrs = {
         "test_binary": attr.label(
             mandatory = True,
-            allow_single_file = True,
+            executable = True,
+            cfg = "target",
         ),
         "data": attr.label_list(allow_empty = True, allow_files = True),
     },
diff --git a/build_tools/bazel/third_party_import/llvm-project/overlay/mlir/BUILD.bazel b/build_tools/bazel/third_party_import/llvm-project/overlay/mlir/BUILD.bazel
index 4aeb2b8..d129c47 100644
--- a/build_tools/bazel/third_party_import/llvm-project/overlay/mlir/BUILD.bazel
+++ b/build_tools/bazel/third_party_import/llvm-project/overlay/mlir/BUILD.bazel
@@ -458,6 +458,7 @@
     srcs = [
         "include/mlir/Dialect/StandardOps/IR/Ops.td",
         "include/mlir/IR/OpAsmInterface.td",
+        "include/mlir/IR/SymbolInterfaces.td",
         "include/mlir/Interfaces/CallInterfaces.td",
         "include/mlir/Interfaces/ControlFlowInterfaces.td",
         "include/mlir/Interfaces/SideEffectInterfaces.td",
diff --git a/build_tools/cmake/iree_copts.cmake b/build_tools/cmake/iree_copts.cmake
index 9d99d9b..dd3b3c0 100644
--- a/build_tools/cmake/iree_copts.cmake
+++ b/build_tools/cmake/iree_copts.cmake
@@ -38,10 +38,6 @@
   ${PROJECT_BINARY_DIR}
 )
 
-if(${IREE_ENABLE_RUNTIME_TRACING})
-  set (CMAKE_EXE_LINKER_FLAGS -ldl)
-endif()
-
 iree_select_compiler_opts(IREE_DEFAULT_COPTS
   CLANG
     # LINT.IfChange(clang_diagnostics)
diff --git a/docs/developing_iree/e2e_benchmarking.md b/docs/developing_iree/e2e_benchmarking.md
index 3fe8604..96523c4 100644
--- a/docs/developing_iree/e2e_benchmarking.md
+++ b/docs/developing_iree/e2e_benchmarking.md
@@ -178,12 +178,19 @@
 # Enter the TensorFlow Bazel workspace.
 cd third_party/tensorflow/
 
-# Build the benchmark_model binary without RUY...
+# Build the benchmark_model binary.
 bazel build --copt=-mavx2 -c opt \
   //tensorflow/lite/tools/benchmark:benchmark_model
 
-# ...or build the benchmark_model binary with RUY. This will overwrite the
+# By default, TFLite/x86 uses various matrix multiplication libraries.
+# It is possible to force it to only use Ruy for all matrix multiplications.
+# That is the default on ARM but not on x86. This will overwrite the
 # previous binary unless you move it.
+#
+# Note that Ruy takes care of -mavx2 and other AVX extensions internally,
+# so this passing this flag here isn't going to make a difference to
+# matrix multiplications. However, the rest of TFLite's kernels outside
+# of ruy will still benefit from -mavx2.
 bazel build --copt=-mavx2 -c opt \
   --define=tflite_with_ruy=true \
   //tensorflow/lite/tools/benchmark:benchmark_model
@@ -274,6 +281,9 @@
 
 ```shell
 # Build the benchmark_model binary without any add-ons.
+# Note that unlike TFLite/x86, TFLite/ARM uses Ruy by default for all
+# matrix multiplications (No need to pass tflite_with_ruy), except for some
+# matrix*vector products. Below we show how to force using ruy also for that.
 bazel build -c opt \
   --config=android_arm64 \
   --cxxopt='--std=c++17' \
@@ -286,20 +296,21 @@
 ```
 
 ```shell
-# Build the benchmark_model binary with ruy.
-bazel build --copt=-mavx2 -c opt \
+# Build the benchmark_model binary using ruy even for matrix*vector
+# products. This is only worth trying in models that are heavy on matrix*vector
+# shapes, typically LSTMs and other RNNs.
+bazel build -c opt \
   --config=android_arm64 \
   --cxxopt='--std=c++17' \
-  --define=tflite_with_ruy=true \
   --copt=-DTFLITE_WITH_RUY_GEMV \
   //tensorflow/lite/tools/benchmark:benchmark_model
 
 # Rename the binary for comparison with the standard benchmark_model.
 mv bazel-bin/tensorflow/lite/tools/benchmark/benchmark_model \
-  bazel-bin/tensorflow/lite/tools/benchmark/benchmark_model_plus_ruy
-adb push bazel-bin/tensorflow/lite/tools/benchmark/benchmark_model_plus_ruy \
+  bazel-bin/tensorflow/lite/tools/benchmark/benchmark_model_plus_ruy_gemv
+adb push bazel-bin/tensorflow/lite/tools/benchmark/benchmark_model_plus_ruy_gemv \
   /data/local/tmp/
-adb shell chmod +x /data/local/tmp/benchmark_model_plus_ruy
+adb shell chmod +x /data/local/tmp/benchmark_model_plus_ruy_gemv
 ```
 
 ```shell
@@ -336,17 +347,15 @@
   --warmup_runs=1 \
   --num_threads=1 \
   --num_runs=10 \
-  --enable_op_profiling=true
 ```
 
 ```shell
-# Benchmark with TFLite + RUY.
-adb shell taskset f0 /data/local/tmp/benchmark_model_plus_ruy \
+# Benchmark with TFLite + RUY GEMV
+adb shell taskset f0 /data/local/tmp/benchmark_model_plus_ruy_gemv \
   --graph=/data/local/tmp/MatrixOpsStaticModule/tflite/matmul_lhs_batch.tflite \
   --warmup_runs=1 \
   --num_threads=1 \
   --num_runs=10 \
-  --enable_op_profiling=true
 ```
 
 ```shell
@@ -356,7 +365,6 @@
   --warmup_runs=1 \
   --num_threads=1 \
   --num_runs=10 \
-  --enable_op_profiling=true
 ```
 
 ```shell
@@ -366,7 +374,6 @@
   --warmup_runs=1 \
   --num_threads=1 \
   --num_runs=10 \
-  --enable_op_profiling=true \
   --use_gpu=true
 ```
 
@@ -384,3 +391,30 @@
 name of the `.tflite` graph that you need to benchmark _may_ be different from
 the name of the trace that you want to benchmark, but you can use `cat` on
 the `graph_path` file to verify the correct `.tflite` filename if you're unsure.
+
+### Profile
+
+There are 2 profilers built into TFLite's `benchmark_model` program. Both of them impact latencies, so they should only be used to get a breakdown of the relative time spent in each operator type, they should not be enabled for the purpose of measuring a latency.
+
+The first is `enable_op_profiling`. It's based on timestamps before and after each op. It's a runtime commandline flag taken by `benchmark_model`. Example:
+
+```
+adb shell taskset f0 /data/local/tmp/benchmark_model \
+  --graph=/data/local/tmp/MatrixOpsStaticModule/tflite/matmul_lhs_batch.tflite \
+  --warmup_runs=1 \
+  --num_threads=1 \
+  --num_runs=10 \
+  --enable_op_profiling=true
+```
+
+The second is `ruy_profiler`. Despite its name, it's available regardless of whether `ruy` is used for the matrix multiplications. It's a sampling profiler, which allows it to provide some more detailed informations, particularly on matrix multiplications. It's a build-time switch:
+
+```
+blaze build \
+  --define=ruy_profiler=true \
+  -c opt \
+  --config=android_arm64 \
+  //tensorflow/lite/tools/benchmark:benchmark_model
+```
+
+The binary thus built can be run like above, no commandline flag needed.
diff --git a/docs/get_started/getting_started_android_cmake.md b/docs/get_started/getting_started_android_cmake.md
index ba78b21..e740b27 100644
--- a/docs/get_started/getting_started_android_cmake.md
+++ b/docs/get_started/getting_started_android_cmake.md
@@ -263,16 +263,16 @@
 ### Dylib LLVM AOT backend
 
 To compile IREE module for the target Android device (assume Android 10 AArc64)
-we need install the corresponding standalone toolchain and setting AOT linker
-path environment variable:
+we need to use the corresponding standalone toolchain (which can be found in
+ANDROID_NDK) and setting AOT linker path environment variable:
 
 ```shell
-$ export ANDROID_ARM64_TOOLCHAIN=/path/to/install/the/toolchain
-$ "${ANDROID_NDK?}/build/tools/make-standalone-toolchain.sh" --arch=arm64 --platform=android-29 \
-    --install-dir="${ANDROID_ARM64_TOOLCHAIN?}"
-$ export IREE_LLVMAOT_LINKER_PATH="${ANDROID_ARM64_TOOLCHAIN?}/aarch64-linux-android/bin/ld"
+$ export IREE_LLVMAOT_LINKER_PATH="${ANDROID_NDK?}/toolchains/llvm/prebuilt/linux-x86_64/bin/aarch64-linux-android29-clang++ -static-libstdc++ -O3"
 ```
 
+`-static-libstdc++` is needed because some dynamic libraries would not be able
+to open.
+
 Translate a source MLIR into an IREE module:
 
 ```shell
diff --git a/iree/base/BUILD b/iree/base/BUILD
index 2d91662..a1b46ba 100644
--- a/iree/base/BUILD
+++ b/iree/base/BUILD
@@ -395,7 +395,12 @@
     name = "synchronization",
     srcs = ["synchronization.c"],
     hdrs = ["synchronization.h"],
-    linkopts = ["-lpthread"],
+    linkopts = select({
+        "//iree:iree_is_msvc": [],
+        "//conditions:default": [
+            "-lpthread",
+        ],
+    }),
     deps = [
         ":api",
         ":atomics",
@@ -460,10 +465,13 @@
     copts = [
         "-D_GNU_SOURCE=1",
     ],
-    linkopts = [
-        "-ldl",
-        "-lpthread",
-    ],
+    linkopts = select({
+        "//iree:iree_is_msvc": [],
+        "//conditions:default": [
+            "-ldl",
+            "-lpthread",
+        ],
+    }),
     deps = [
         ":api",
         ":atomics",
diff --git a/iree/base/CMakeLists.txt b/iree/base/CMakeLists.txt
index 4e6b7ce..257618b 100644
--- a/iree/base/CMakeLists.txt
+++ b/iree/base/CMakeLists.txt
@@ -595,6 +595,11 @@
     iree::testing::gtest_main
 )
 
+iree_select_compiler_opts(IREE_LINKOPTS_TRACING
+  GCC_OR_CLANG
+    -ldl
+)
+
 if(${IREE_ENABLE_RUNTIME_TRACING})
   iree_cc_library(
     NAME
@@ -605,6 +610,8 @@
       "${IREE_ROOT_DIR}/third_party/tracy/TracyC.h"
     SRCS
       "tracing.cc"
+    LINKOPTS
+      ${IREE_LINKOPTS_TRACING}
     DEPS
       absl::core_headers
     DEFINES
diff --git a/iree/build_defs.oss.bzl b/iree/build_defs.oss.bzl
index 5364f6b..4d359e7 100644
--- a/iree/build_defs.oss.bzl
+++ b/iree/build_defs.oss.bzl
@@ -85,11 +85,14 @@
 def cc_binary(linkopts = [], **kwargs):
     """Wrapper around low-level cc_binary that adds flags."""
     _cc_binary(
-        linkopts = linkopts + [
-            # Just include libraries that should be presumed in 2020.
-            "-ldl",
-            "-lpthread",
-        ],
+        linkopts = linkopts + select({
+            "//iree:iree_is_msvc": [],
+            "//conditions:default": [
+                # Just include libraries that should be presumed in 2020.
+                "-ldl",
+                "-lpthread",
+            ],
+        }),
         **kwargs
     )
 
diff --git a/iree/compiler/Conversion/HLOToLinalg/test/pad.mlir b/iree/compiler/Conversion/HLOToLinalg/test/pad.mlir
index a42a079..29c700b 100644
--- a/iree/compiler/Conversion/HLOToLinalg/test/pad.mlir
+++ b/iree/compiler/Conversion/HLOToLinalg/test/pad.mlir
@@ -76,14 +76,10 @@
     hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write"
   }
 }
-// TODO(hanchung): Make it just a copy op.
 // CHECK_LABEL: @pad_no_op
-//   CHECK-DAG: %[[CST:.+]] = constant 0.000000e+00 : f32
 //   CHECK-DAG: %[[OUT:.+]] = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<12x4xf32>
 //   CHECK-DAG: %[[IN:.+]] = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<12x4xf32>
-//       CHECK: linalg.fill(%[[OUT]], %[[CST]])
-//       CHECK: %[[SUBVIEW:.+]] = subview %[[OUT]][0, 0] [12, 4] [1, 1]
-//       CHECK: linalg.copy(%[[IN]], %[[SUBVIEW]])
+//   CHECK: linalg.copy(%[[IN]], %[[OUT]])
 
 // -----
 
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/ConvertToSPIRVPass.cpp b/iree/compiler/Conversion/LinalgToSPIRV/ConvertToSPIRVPass.cpp
index 899a485..93b8809 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/ConvertToSPIRVPass.cpp
+++ b/iree/compiler/Conversion/LinalgToSPIRV/ConvertToSPIRVPass.cpp
@@ -204,7 +204,7 @@
     if (vecType.getRank() != 2) return failure();
     // TODO(thomasraoux): use coloumn major operand when TransfertRead +
     // TransposeOp.
-    if (!op.permutation_map().isIdentity()) return failure();
+    if (!op.permutation_map().isMinorIdentity()) return failure();
     if (op.masked() &&
         llvm::any_of(op.masked()->template cast<ArrayAttr>(),
                      [](mlir::Attribute maskedDim) {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_spirv.mlir b/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_spirv.mlir
index 52aeff7..ac83d70 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_spirv.mlir
+++ b/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_spirv.mlir
@@ -104,3 +104,41 @@
     return
   }
 }
+
+// -----
+
+#map0 = affine_map<(d0, d1) -> (d1)>
+#map1 = affine_map<(d0, d1, d2) -> (d0, d2)>
+#map2 = affine_map<(d0, d1, d2) -> (d2, d1)>
+#map3 = affine_map<(d0, d1, d2) -> (d0, d1)>
+
+module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader, CooperativeMatrixNV, Int8, Float16, StorageUniform16, StorageBuffer8BitAccess, Float16Buffer], [SPV_KHR_storage_buffer_storage_class, SPV_NV_cooperative_matrix, SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
+  func @kernel_matmul_vector_memref(%arg0: memref<4096x256xvector<4xi32>>, %arg1: memref<4096x256xvector<4xi32>>, %arg2: memref<4096x1024xvector<4xi32>>) attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
+    %c32 = constant 32 : index
+    %c4096 = constant 4096 : index
+    %c0 = constant 0 : index
+    %cst = constant dense<0> : vector<4xi32>
+    // CHECK: %[[C:.+]] = spv.CooperativeMatrixLoadNV %{{.*}}, %{{.*}}, %{{.*}}
+    %4 = vector.transfer_read %arg2[%c0, %c0], %cst : memref<4096x1024xvector<4xi32>>, vector<16x16xi32>
+    // CHECK: %[[ACC:.+]] = spv.Variable : !spv.ptr<!spv.coopmatrix<16x16xi32, Subgroup>, Function>
+    // CHECK: spv.loop {
+      // CHECK: spv.Branch ^[[BB:.+]](%{{.*}}, %[[C]] : i32, !spv.coopmatrix<16x16xi32, Subgroup>)
+      // CHECK: ^[[BB]](%{{.*}}: i32, %[[C1:.+]]: !spv.coopmatrix<16x16xi32, Subgroup>)
+    %5 = scf.for %arg3 = %c0 to %c4096 step %c32 iter_args(%arg4 = %4) -> (vector<16x16xi32>) {
+      // CHECK: %[[A:.+]] = spv.CooperativeMatrixLoadNV %{{.*}}, %{{.*}}, %{{.*}}
+      %6 = vector.transfer_read %arg0[%c0, %arg3], %cst : memref<4096x256xvector<4xi32>>, vector<16x32xi8>
+      // CHECK: %[[B:.+]] = spv.CooperativeMatrixLoadNV %{{.*}}, %{{.*}}, %{{.*}}
+      %7 = vector.transfer_read %arg1[%arg3, %c0], %cst : memref<4096x256xvector<4xi32>>, vector<32x16xi8>
+      // CHECK: %[[R:.+]] = spv.CooperativeMatrixMulAddNV %[[A]], %[[B]], %[[C1]]
+      %8 = vector.contract {indexing_maps = [#map1, #map2, #map3], iterator_types = ["parallel", "parallel", "reduction"]} %6, %7, %arg4 : vector<16x32xi8>, vector<32x16xi8> into vector<16x16xi32>
+      // CHECK: spv.Store "Function" %[[ACC]], %[[R]] : !spv.coopmatrix<16x16xi32, Subgroup>
+      // CHECK: spv.Branch ^[[BB]](%{{.*}}, %[[R]] : i32, !spv.coopmatrix<16x16xi32, Subgroup>)
+      scf.yield %8 : vector<16x16xi32>
+    }
+    // CHECK: %[[ACCv:.+]] = spv.Load "Function" %[[ACC]] : !spv.coopmatrix<16x16xi32, Subgroup>
+    // CHECK: spv.CooperativeMatrixStoreNV %{{.*}}, %[[ACCv]], %{{.*}}, %{{.*}}
+    vector.transfer_write %5, %arg2[%c0, %c0] : vector<16x16xi32>, memref<4096x1024xvector<4xi32>>
+    return
+  }
+}
+
diff --git a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/SPIRVToMSL.cpp b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/SPIRVToMSL.cpp
index 2f5abb7..5d89aa3 100644
--- a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/SPIRVToMSL.cpp
+++ b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/SPIRVToMSL.cpp
@@ -14,6 +14,9 @@
 
 #include "iree/compiler/Dialect/HAL/Target/MetalSPIRV/SPIRVToMSL.h"
 
+#include <vector>
+
+#include "llvm/ADT/STLExtras.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/raw_ostream.h"
 #include "spirv_cross/spirv_msl.hpp"
@@ -37,6 +40,66 @@
     if (workgroupSize.constant != 0) return {0, 0, 0};
     return {workgroupSize.x, workgroupSize.y, workgroupSize.z};
   }
+
+  // A struct containing a resource descriptor's information.
+  struct Descriptor {
+    uint32_t set;
+    uint32_t binding;
+
+    Descriptor(uint32_t s, uint32_t b) : set(s), binding(b) {}
+
+    friend bool operator<(const Descriptor& l, const Descriptor& r) {
+      return std::tie(l.set, l.binding) < std::tie(r.set, r.binding);
+    }
+  };
+
+  // Returns all all resource buffer descriptors' set and binding number pairs
+  // in increasing order.
+  std::vector<Descriptor> getBufferSetBindingPairs() {
+    std::vector<Descriptor> descriptors;
+
+    // Iterate over all variables in the SPIR-V blob.
+    ir.for_each_typed_id<SPIRV_CROSS_NAMESPACE::SPIRVariable>(
+        [&](uint32_t id, SPIRV_CROSS_NAMESPACE::SPIRVariable& var) {
+          auto storage = var.storage;
+          switch (storage) {
+              // Non-interface variables. We don't care.
+            case spv::StorageClassFunction:
+            case spv::StorageClassPrivate:
+            case spv::StorageClassWorkgroup:
+              // Builtin variables. We don't care either.
+            case spv::StorageClassInput:
+              return;
+            default:
+              break;
+          }
+          if (storage == spv::StorageClassUniform ||
+              storage == spv::StorageClassStorageBuffer) {
+            uint32_t setNo = get_decoration(id, spv::DecorationDescriptorSet);
+            uint32_t bindingNo = get_decoration(id, spv::DecorationBinding);
+            descriptors.emplace_back(setNo, bindingNo);
+            return;
+          }
+          // TODO(antiagainst): push constant
+          llvm_unreachable("unspported storage class in SPIRVToMSLCompiler");
+        });
+
+    llvm::sort(descriptors);
+    return descriptors;
+  }
+
+  Options getCompilationOptions() {
+    // TODO(antiagainst): fill out the following according to the Metal GPU
+    // family.
+    SPIRVToMSLCompiler::Options spvCrossOptions;
+    spvCrossOptions.platform = SPIRVToMSLCompiler::Options::Platform::macOS;
+    spvCrossOptions.msl_version =
+        SPIRVToMSLCompiler::Options::make_msl_version(2, 0);
+    // Eanble using Metal argument buffers. It is more akin to Vulkan descriptor
+    // sets, which is how IREE HAL models resource bindings and mappings.
+    spvCrossOptions.argument_buffers = true;
+    return spvCrossOptions;
+  }
 };
 }  // namespace
 
@@ -49,15 +112,25 @@
   spvCrossCompiler.set_entry_point(
       entryPoint, spv::ExecutionModel::ExecutionModelGLCompute);
 
-  // TODO(antiagainst): fill out the following according to the Metal GPU
-  // family.
-  SPIRVToMSLCompiler::Options spvCrossOptions;
-  spvCrossOptions.platform = SPIRVToMSLCompiler::Options::Platform::macOS;
-  spvCrossOptions.msl_version =
-      SPIRVToMSLCompiler::Options::make_msl_version(2, 0);
-  // Eanble using Metal argument buffers. It is more akin to Vulkan descriptor
-  // sets, which is how IREE HAL models resource bindings and mappings.
-  spvCrossOptions.argument_buffers = true;
+  // Explicitly set the argument buffer index for each SPIR-V resource variable.
+  auto descriptors = spvCrossCompiler.getBufferSetBindingPairs();
+  for (const auto& descriptor : descriptors) {
+    if (descriptor.set != 0) {
+      llvm_unreachable(
+          "multiple descriptor set unimplemented in SPIRVToMSLCompiler");
+    }
+
+    SPIRV_CROSS_NAMESPACE::MSLResourceBinding binding = {};
+    binding.stage = spv::ExecutionModelGLCompute;
+    binding.desc_set = descriptor.set;
+    binding.binding = descriptor.binding;
+    // We only interact with buffers in IREE.
+    binding.msl_buffer = descriptor.binding;
+
+    spvCrossCompiler.add_msl_resource_binding(binding);
+  }
+
+  auto spvCrossOptions = spvCrossCompiler.getCompilationOptions();
   spvCrossCompiler.set_msl_options(spvCrossOptions);
 
   std::string mslSource = spvCrossCompiler.compile();
diff --git a/iree/test/e2e/xla_ops/BUILD b/iree/test/e2e/xla_ops/BUILD
index f60833f..a3c9060 100644
--- a/iree/test/e2e/xla_ops/BUILD
+++ b/iree/test/e2e/xla_ops/BUILD
@@ -37,24 +37,34 @@
     "abs.mlir"
     "add.mlir"
     "broadcast.mlir"
+    "broadcast_add.mlir"
     "broadcast_in_dim.mlir"
-    "convert.mlir"
+    "clamp.mlir"
+    "compare.mlir"
     "constant.mlir"
+    "convert.mlir"
     "cosine.mlir"
+    "divide.mlir"
     "exponential.mlir"
+    "gather.mlir"
     "log.mlir"
     "log_plus_one.mlir"
     "maximum.mlir"
     "minimum.mlir"
     "multiply.mlir"
     "negate.mlir"
+    "remainder.mlir"
     "reshape.mlir"
     "rsqrt.mlir"
+    "select.mlir"
     "sine.mlir"
     "slice.mlir"
     "sqrt.mlir"
+    "subtract.mlir"
     "tanh.mlir"
+    "torch_index_select.mlir"
     "transpose.mlir"
+    "while.mlir"
   TARGET_BACKEND
     metal-spirv
   DRIVER
diff --git a/iree/test/e2e/xla_ops/CMakeLists.txt b/iree/test/e2e/xla_ops/CMakeLists.txt
index 7d8a568..1f89648 100644
--- a/iree/test/e2e/xla_ops/CMakeLists.txt
+++ b/iree/test/e2e/xla_ops/CMakeLists.txt
@@ -19,24 +19,34 @@
     "abs.mlir"
     "add.mlir"
     "broadcast.mlir"
+    "broadcast_add.mlir"
     "broadcast_in_dim.mlir"
-    "convert.mlir"
+    "clamp.mlir"
+    "compare.mlir"
     "constant.mlir"
+    "convert.mlir"
     "cosine.mlir"
+    "divide.mlir"
     "exponential.mlir"
+    "gather.mlir"
     "log.mlir"
     "log_plus_one.mlir"
     "maximum.mlir"
     "minimum.mlir"
     "multiply.mlir"
     "negate.mlir"
+    "remainder.mlir"
     "reshape.mlir"
     "rsqrt.mlir"
+    "select.mlir"
     "sine.mlir"
     "slice.mlir"
     "sqrt.mlir"
+    "subtract.mlir"
     "tanh.mlir"
+    "torch_index_select.mlir"
     "transpose.mlir"
+    "while.mlir"
   TARGET_BACKEND
     metal-spirv
   DRIVER
diff --git a/scripts/prepare_doc_publication.py b/scripts/prepare_doc_publication.py
index 716e31a..352bcf0 100755
--- a/scripts/prepare_doc_publication.py
+++ b/scripts/prepare_doc_publication.py
@@ -86,7 +86,7 @@
     # 'Getting Started' is 3.
     # 'Developing IREE' is 4.
     'design_roadmap.md': 5,
-    'milestones.md': 6,
+    'objectives.md': 6,
     'xla_op_coverage.md': 7,
     'tf_e2e_coverage.md': 8,
     'iree_community.md': 9,
diff --git a/third_party/llvm-project b/third_party/llvm-project
index 51ff045..f402e68 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit 51ff04567b2f8d06b2062bd3ed72eab2e93e4466
+Subproject commit f402e682d0ef5598eeffc9a21a691b03e602ff58
diff --git a/third_party/tensorflow b/third_party/tensorflow
index 74747f0..2d6bdab 160000
--- a/third_party/tensorflow
+++ b/third_party/tensorflow
@@ -1 +1 @@
-Subproject commit 74747f0c6b017df8fe87d1912f8eb4e2e287fbda
+Subproject commit 2d6bdab3adb0b8949763d5c63426338f938c9efe