diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS
index 56d23b8..0aac245 100644
--- a/.github/CODEOWNERS
+++ b/.github/CODEOWNERS
@@ -21,6 +21,7 @@
 
 # Bindings
 /bindings/python/ @stellaraccident
+/runtime/bindings/python/ @stellaraccident
 /bindings/tflite/ @benvanik
 
 # Integrations
@@ -41,9 +42,9 @@
 /build_tools/ @GMNGeoffrey @ScottTodd
 /build_tools/benchmarks/ @GMNGeoffrey @antiagainst
 /build_tools/python_deploy/ @stellaraccident
+/build_tools/scripts/ @GMNGeoffrey @ScottTodd
 /build_tools/third_party/ @GMNGeoffrey @ScottTodd @stellaraccident
 /.github/ @GMNGeoffrey @ScottTodd
-/scripts/ @GMNGeoffrey @ScottTodd
 
 # llvm-external-projects
 /llvm-external-projects/ @stellaraccident
@@ -52,7 +53,6 @@
 /llvm-external-projects/iree-dialects/test/iree_linalgext @hanhanW @MaheshRavishankar
 
 # Other Top-Level Directories
-/colab/ @ScottTodd
 /docs/ @ScottTodd
 
 
@@ -74,11 +74,12 @@
 
 
 # Runtime
-/iree/hal/cuda/ @ThomasRaoux
-/iree/hal/cts/ @ScottTodd
-/iree/hal/vulkan/ @antiagainst @ScottTodd
+/runtime/src/iree/ @benvanik
+/runtime/src/iree/hal/cuda/ @ThomasRaoux
+/runtime/src/iree/hal/cts/ @ScottTodd
+/runtime/src/iree/hal/vulkan/ @antiagainst @ScottTodd
 
 # Other IREE directories
-/iree/samples/ @ScottTodd
+/samples/ @ScottTodd
 /iree/test/ @ghost
 /iree/tools/ @benvanik @GMNGeoffrey
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 5921cd7..9828641 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -46,6 +46,7 @@
 option(IREE_BUILD_DOCS "Builds IREE docs." OFF)
 option(IREE_BUILD_SAMPLES "Builds IREE sample projects." ON)
 option(IREE_BUILD_TRACY "Builds tracy server tools." OFF)
+option(IREE_BUILD_TORCH_MLIR_SUPPORT "Builds support for compiling torch-mlir programs." ON)
 
 option(IREE_BYTECODE_MODULE_FORCE_SYSTEM_DYLIB_LINKER "Use the system linker when generating IREE modules in tests/samples/benchmarks (useful for Tracy)." OFF)
 
@@ -67,7 +68,6 @@
 option(IREE_BUILD_EXPERIMENTAL_REMOTING "Builds experimental remoting support." OFF)
 option(IREE_BUILD_EXPERIMENTAL_WEB_SAMPLES "Builds experimental web samples." OFF)
 option(IREE_HAL_DRIVER_EXPERIMENTAL_ROCM "Builds the experimental ROCm Backend." OFF)
-option(IREE_BUILD_TORCH_MLIR_SUPPORT "Builds support for compiling torch-mlir programs." ON)
 
 #-------------------------------------------------------------------------------
 # Derived flags based on primary options
diff --git a/build_tools/cmake/build_runtime_emscripten.sh b/build_tools/cmake/build_runtime_emscripten.sh
index 9cbb7da..842326c 100755
--- a/build_tools/cmake/build_runtime_emscripten.sh
+++ b/build_tools/cmake/build_runtime_emscripten.sh
@@ -48,4 +48,4 @@
   -DIREE_BUILD_SAMPLES=ON
 
 # TODO(scotttodd): expand this list of targets
-"${CMAKE_BIN?}" --build . --target iree_samples_simple_embedding_simple_embedding_vmvx_sync -- -k 0
+"${CMAKE_BIN?}" --build . --target iree_samples_iree_simple_embedding_simple_embedding_vmvx_sync -- -k 0
diff --git a/build_tools/cmake/iree_bytecode_module.cmake b/build_tools/cmake/iree_bytecode_module.cmake
index 24ecae2..ed72ede 100644
--- a/build_tools/cmake/iree_bytecode_module.cmake
+++ b/build_tools/cmake/iree_bytecode_module.cmake
@@ -109,7 +109,7 @@
       ${_RULE_SRC}
       ${_RULE_DEPENDS}
     COMMENT
-      "Generating VMFB for ${_FRIENDLY_NAME}"
+      "Generating ${_MODULE_FILE_NAME} from ${_FRIENDLY_NAME}"
     VERBATIM
   )
 
diff --git a/build_tools/testing/test_samples.sh b/build_tools/testing/test_samples.sh
index b7da1e4..1387a6d 100755
--- a/build_tools/testing/test_samples.sh
+++ b/build_tools/testing/test_samples.sh
@@ -14,5 +14,5 @@
 git submodule sync
 git submodule update --init
 
-./iree/samples/dynamic_shapes/test.sh
-./iree/samples/variables_and_state/test.sh
+./samples/dynamic_shapes/test.sh
+./samples/variables_and_state/test.sh
diff --git a/docs/developers/developing_iree/developer_overview.md b/docs/developers/developing_iree/developer_overview.md
index 0866fb8..36e2596 100644
--- a/docs/developers/developing_iree/developer_overview.md
+++ b/docs/developers/developing_iree/developer_overview.md
@@ -3,6 +3,9 @@
 This guide provides an overview of IREE's project structure and main tools for
 developers.
 
+** Note: project layout is evolving at the moment, see
+   https://github.com/google/iree/issues/8955 **
+
 ## Project Code Layout
 
 [iree/](https://github.com/google/iree/blob/main/iree/)
@@ -13,31 +16,38 @@
 
 *   Integrations between IREE and other frameworks, such as TensorFlow
 
+[runtime/](https://github.com/google/iree/tree/main/runtime/)
+
+*   IREE runtime code, with no dependencies on the compiler
+
 [bindings/](https://github.com/google/iree/blob/main/bindings/)
 
 *   Language and platform bindings, such as Python
+*   Also see [runtime/bindings/](https://github.com/google/iree/tree/main/runtime/bindings)
 
-[colab/](https://github.com/google/iree/blob/main/colab/)
+[samples/](https://github.com/google/iree/blob/main/samples/)
 
-*   Colab notebooks for interactively using IREE's Python bindings
+*   Samples built using IREE's runtime and compiler
+*   Also see the separate https://github.com/google/iree-samples repository
 
-## IREE Code Layout
-
-[iree/base/](https://github.com/google/iree/blob/main/iree/base/)
-
-*   Common types and utilities used throughout IREE
+## IREE Compiler Code Layout
 
 [iree/compiler/](https://github.com/google/iree/blob/main/iree/compiler/)
 
 *   IREE's MLIR dialects, LLVM compiler passes, module translation code, etc.
-    Code here should not depend on anything in the runtime
 
-[iree/hal/](https://github.com/google/iree/blob/main/iree/hal/)
+## IREE Runtime Code Layout
+
+[iree/base/](https://github.com/google/iree/blob/main/runtime/src/iree/base/)
+
+*   Common types and utilities used throughout the runtime
+
+[iree/hal/](https://github.com/google/iree/blob/main/runtime/src/iree/hal/)
 
 *   **H**ardware **A**bstraction **L**ayer for IREE's runtime, with
     implementations for hardware and software backends
 
-[iree/schemas/](https://github.com/google/iree/blob/main/iree/schemas/)
+[iree/schemas/](https://github.com/google/iree/blob/main/runtime/src/iree/schemas/)
 
 *   Shared data storage format definitions, primarily using
     [FlatBuffers](https://google.github.io/flatbuffers/)
@@ -46,7 +56,7 @@
 
 *   Assorted tools used to optimize, translate, and evaluate IREE
 
-[iree/vm/](https://github.com/google/iree/blob/main/iree/vm/)
+[iree/vm/](https://github.com/google/iree/blob/main/runtime/src/iree/vm/)
 
 *   Bytecode **V**irtual **M**achine used to work with IREE modules and invoke
     IREE functions
diff --git a/docs/developers/get_started/building_with_emscripten.md b/docs/developers/get_started/building_with_emscripten.md
index e97f04d..35b4eaf 100644
--- a/docs/developers/get_started/building_with_emscripten.md
+++ b/docs/developers/get_started/building_with_emscripten.md
@@ -55,7 +55,7 @@
 
 ```
 cmake --build ../iree-build-emscripten/ \
-  --target iree_samples_simple_embedding_simple_embedding_vmvx_sync
+  --target iree_samples_iree_simple_embedding_simple_embedding_vmvx_sync
 ```
 
 ### Load into a WebAssembly Environment
diff --git a/docs/website/docs/bindings/c-api.md b/docs/website/docs/bindings/c-api.md
index e140b22..4de69e9 100644
--- a/docs/website/docs/bindings/c-api.md
+++ b/docs/website/docs/bindings/c-api.md
@@ -7,11 +7,11 @@
 
 | Component header file                                                       | Overview                                                                  |
 |-----------------------------------------------------------------------------|---------------------------------------------------------------------------|
-| [iree/base/api.h](https://github.com/google/iree/blob/main/iree/base/api.h) | Core API, type definitions, ownership policies, utilities                 |
-| [iree/vm/api.h](https://github.com/google/iree/blob/main/iree/vm/api.h)     | VM APIs: loading modules, I/O, calling functions                          |
-| [iree/hal/api.h](https://github.com/google/iree/blob/main/iree/hal/api.h)   | HAL APIs: device management, synchronization, accessing hardware features |
+| [iree/base/api.h](https://github.com/google/iree/blob/main/runtime/src/iree/base/api.h) | Core API, type definitions, ownership policies, utilities                 |
+| [iree/vm/api.h](https://github.com/google/iree/blob/main/runtime/src/iree/vm/api.h)     | VM APIs: loading modules, I/O, calling functions                          |
+| [iree/hal/api.h](https://github.com/google/iree/blob/main/runtime/src/iree/hal/api.h)   | HAL APIs: device management, synchronization, accessing hardware features |
 
-The [samples/](https://github.com/google/iree/tree/main/iree/samples)
+The [samples/](https://github.com/google/iree/tree/main/samples)
 directory demonstrates several ways to use IREE's C API.
 
 ## Prerequisites
diff --git a/docs/website/docs/bindings/python.md b/docs/website/docs/bindings/python.md
index 042cc98..9fc16b7 100644
--- a/docs/website/docs/bindings/python.md
+++ b/docs/website/docs/bindings/python.md
@@ -115,9 +115,9 @@
 [readthedocs](https://iree-python-api.readthedocs.io/en/latest/).
 
 Check out the samples in IREE's
-[colab/ directory](https://github.com/google/iree/tree/main/colab) and the
-[iree-samples repository](https://github.com/google/iree-samples) for examples
-using the Python APIs.
+[samples/colab/ directory](https://github.com/google/iree/tree/main/samples/colab)
+and the [iree-samples repository](https://github.com/google/iree-samples) for
+examples using the Python APIs.
 
 <!-- ## Troubleshooting -->
 
diff --git a/docs/website/docs/blog/2021-07-19-tflite-tosa.md b/docs/website/docs/blog/2021-07-19-tflite-tosa.md
index 9da0139..07e83f7 100644
--- a/docs/website/docs/blog/2021-07-19-tflite-tosa.md
+++ b/docs/website/docs/blog/2021-07-19-tflite-tosa.md
@@ -34,7 +34,7 @@
 ## Examples
 
 TFLite with IREE is available in Python and Java.  We have a
-[colab notebook](https://colab.research.google.com/github/google/iree/blob/main/colab/tflite_text_classification.ipynb)
+[colab notebook](https://colab.research.google.com/github/google/iree/blob/main/samples/colab/tflite_text_classification.ipynb)
 that shows how to use IREE’s python bindings and TFLite compiler tools to
 compile a pre-trained TFLite model from a flatbuffer and run using IREE.  We
 also have an
diff --git a/docs/website/docs/building-from-source/python-bindings-and-importers.md b/docs/website/docs/building-from-source/python-bindings-and-importers.md
index da04d84..de004f0 100644
--- a/docs/website/docs/building-from-source/python-bindings-and-importers.md
+++ b/docs/website/docs/building-from-source/python-bindings-and-importers.md
@@ -34,8 +34,8 @@
 
 * **`IREE_BUILD_PYTHON_BINDINGS`** : `BOOL`
 
-    Enables building of Python bindings under `bindings/python` in the repository.
-    Defaults to `OFF`.
+    Enables building of Python bindings under `runtime/bindings/python` in the
+    repository. Defaults to `OFF`.
 
 * **`Python3_EXECUTABLE`** : `PATH`
 
diff --git a/docs/website/docs/getting-started/index.md b/docs/website/docs/getting-started/index.md
index 2429f1e..203f70a 100644
--- a/docs/website/docs/getting-started/index.md
+++ b/docs/website/docs/getting-started/index.md
@@ -27,9 +27,10 @@
 
 ## Samples
 
-Check out the samples in IREE's [colab/directory](https://github.com/google/iree/tree/main/colab),
-as well as the [iree-samples repository](https://github.com/google/iree-samples)
-respository, which contains workflow comparisons across frameworks.
+Check out the samples in IREE's
+[samples/colab/ directory](https://github.com/google/iree/tree/main/colab),
+as well as the [iree-samples repository](https://github.com/google/iree-samples),
+which contains workflow comparisons across frameworks.
 
 ## Import
 
@@ -42,7 +43,7 @@
 * Validate only IREE compatible operations remain
 * Write the remaining IR to a file
 
-This fully legalized form can then be compiled without dependencies on the 
+This fully legalized form can then be compiled without dependencies on the
 source model language.
 
 ## Compilation
diff --git a/docs/website/docs/getting-started/tensorflow.md b/docs/website/docs/getting-started/tensorflow.md
index e112fc4..3774542 100644
--- a/docs/website/docs/getting-started/tensorflow.md
+++ b/docs/website/docs/getting-started/tensorflow.md
@@ -97,10 +97,10 @@
 
 | Colab notebooks |  |
 | -- | -- |
-Training an MNIST digits classifier | [![Open in Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/colab/mnist_training.ipynb)
-Edge detection module | [![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/colab/edge_detection.ipynb)
-Pretrained ResNet50 inference | [![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/colab/resnet.ipynb)
-TensorFlow Hub Import | [![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/colab/tensorflow_hub_import.ipynb)
+Training an MNIST digits classifier | [![Open in Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/colab/mnist_training.ipynb)
+Edge detection module | [![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/colab/edge_detection.ipynb)
+Pretrained ResNet50 inference | [![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/colab/resnet.ipynb)
+TensorFlow Hub Import | [![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/colab/tensorflow_hub_import.ipynb)
 
 End-to-end execution tests can be found in IREE's
 [integrations/tensorflow/e2e/](https://github.com/google/iree/tree/main/integrations/tensorflow/e2e)
diff --git a/docs/website/docs/getting-started/tflite.md b/docs/website/docs/getting-started/tflite.md
index 9f56962..ddb8f01 100644
--- a/docs/website/docs/getting-started/tflite.md
+++ b/docs/website/docs/getting-started/tflite.md
@@ -130,11 +130,11 @@
 
 ## Troubleshooting
 
-Failures during the import step usually indicate a failure to lower from 
+Failures during the import step usually indicate a failure to lower from
 TensorFlow Lite's operations to TOSA, the intermediate representation used by
 IREE. Many TensorFlow Lite operations are not fully supported, particularly
 those than use dynamic shapes. File an issue to IREE's TFLite model support
-[project](https://github.com/google/iree/projects/42). 
+[project](https://github.com/google/iree/projects/42).
 
 
 ## Additional Samples
@@ -152,7 +152,7 @@
 
 | Colab notebooks |  |
 | -- | -- |
-Text classification with TFLite and IREE | [![Open in Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/colab/tflite_text_classification.ipynb)
+Text classification with TFLite and IREE | [![Open in Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/colab/tflite_text_classification.ipynb)
 
 !!! todo
 
@@ -161,4 +161,3 @@
     [Java TFLite bindings](https://github.com/google/iree/tree/main/bindings/tflite/java),
     once it is complete at
     [not-jenni/iree-android-tflite-demo](https://github.com/not-jenni/iree-android-tflite-demo).
-
diff --git a/experimental/web/sample_dynamic/README.md b/experimental/web/sample_dynamic/README.md
index 2e1bec1..cd36f46 100644
--- a/experimental/web/sample_dynamic/README.md
+++ b/experimental/web/sample_dynamic/README.md
@@ -23,7 +23,7 @@
 to compile the runtime into WebAssembly and JavaScript files.
 
 Any supported IREE program, such as
-[simple_abs.mlir](../../../iree/samples/models/simple_abs.mlir), is compiled using
+[simple_abs.mlir](../../../samples/models/simple_abs.mlir), is compiled using
 the "system library" linking mode. This creates a shared object (typically
 .so/.dll, .wasm in this case). When the runtime attempts to load this file
 using `dlopen()` and `dlsym()`, Emscripten makes use of its
diff --git a/experimental/web/sample_dynamic/build_sample.sh b/experimental/web/sample_dynamic/build_sample.sh
index 6fdc93e..534e4c3 100755
--- a/experimental/web/sample_dynamic/build_sample.sh
+++ b/experimental/web/sample_dynamic/build_sample.sh
@@ -57,7 +57,7 @@
 }
 
 echo "=== Compiling sample MLIR files to VM flatbuffer outputs (.vmfb) ==="
-compile_sample "simple_abs"     "${ROOT_DIR?}/iree/samples/models/simple_abs.mlir"
+compile_sample "simple_abs"     "${ROOT_DIR?}/samples/models/simple_abs.mlir"
 compile_sample "fullyconnected" "${ROOT_DIR?}/iree/test/e2e/models/fullyconnected.mlir"
 compile_sample "collatz"        "${ROOT_DIR?}/iree/test/e2e/models/collatz.mlir"
 
diff --git a/experimental/web/sample_dynamic/serve_sample.sh b/experimental/web/sample_dynamic/serve_sample.sh
index fd1d5e1..d83f873 100755
--- a/experimental/web/sample_dynamic/serve_sample.sh
+++ b/experimental/web/sample_dynamic/serve_sample.sh
@@ -11,4 +11,4 @@
 
 echo "=== Running local webserver, open at http://localhost:8000/ ==="
 
-python3 ${ROOT_DIR?}/scripts/local_web_server.py --directory ${BINARY_DIR}
+python3 ${ROOT_DIR?}/build_tools/scripts/local_web_server.py --directory ${BINARY_DIR}
diff --git a/experimental/web/sample_static/CMakeLists.txt b/experimental/web/sample_static/CMakeLists.txt
index c0210f7..2d78f7b 100644
--- a/experimental/web/sample_static/CMakeLists.txt
+++ b/experimental/web/sample_static/CMakeLists.txt
@@ -51,7 +51,7 @@
 
 target_link_options(${_NAME} PRIVATE
   # https://emscripten.org/docs/porting/connecting_cpp_and_javascript/Interacting-with-code.html#interacting-with-code-ccall-cwrap
-  "-sEXPORTED_FUNCTIONS=['_setup_sample', '_cleanup_sample', '_run_sample']"
+  "-sEXPORTED_FUNCTIONS=['_setup_sample', '_cleanup_sample', '_run_sample', '_malloc']"
   "-sEXPORTED_RUNTIME_METHODS=['ccall','cwrap']"
   #
   "-sASSERTIONS=1"
@@ -96,7 +96,7 @@
 
 target_link_options(${_NAME} PRIVATE
   # https://emscripten.org/docs/porting/connecting_cpp_and_javascript/Interacting-with-code.html#interacting-with-code-ccall-cwrap
-  "-sEXPORTED_FUNCTIONS=['_setup_sample', '_cleanup_sample', '_run_sample']"
+  "-sEXPORTED_FUNCTIONS=['_setup_sample', '_cleanup_sample', '_run_sample', '_malloc']"
   "-sEXPORTED_RUNTIME_METHODS=['ccall','cwrap']"
   #
   "-sASSERTIONS=1"
@@ -132,5 +132,7 @@
   # IREE is pretty good about not allocating outside of startup, so concerns
   # about this causing slow access to memory *may* not affect IREE too much.
   # "-sALLOW_MEMORY_GROWTH=1"
+  # TODO(scotttodd): tune this (figure out where memory is going and trim)
+  # "-sINITIAL_MEMORY=33554432"
   # ------------------------------------------------------------------------- #
 )
diff --git a/experimental/web/sample_static/README.md b/experimental/web/sample_static/README.md
index 7360d3b..0375e0f 100644
--- a/experimental/web/sample_static/README.md
+++ b/experimental/web/sample_static/README.md
@@ -22,10 +22,10 @@
 
 ## How it works
 
-This [MNIST model](../../../iree/samples/models/mnist.mlir), also used in the
-[Vision sample](../../../iree/samples/vision/), is compiled using the "static
+This [MNIST model](../../../samples/models/mnist.mlir), also used in the
+[Vision sample](../../../samples/iree_vision_inference/), is compiled using the "static
 library" output setting of IREE's compiler (see the
-[Static library sample](../../../iree/samples/static_library)). The resulting
+[Static library sample](../../../samples/iree_static_library)). The resulting
 `.h` and `.o` files are compiled together with `main.c`, while the `.vmfb` is
 embedded into a C file that is similarly linked in.
 
diff --git a/experimental/web/sample_static/build_sample.sh b/experimental/web/sample_static/build_sample.sh
index db2cc03..e6b4c0f 100755
--- a/experimental/web/sample_static/build_sample.sh
+++ b/experimental/web/sample_static/build_sample.sh
@@ -46,7 +46,7 @@
 COMPILE_TOOL="${INSTALL_ROOT?}/bin/iree-compile"
 EMBED_DATA_TOOL="${INSTALL_ROOT?}/bin/generate_embed_data"
 INPUT_NAME="mnist"
-INPUT_PATH="${ROOT_DIR?}/iree/samples/models/mnist.mlir"
+INPUT_PATH="${ROOT_DIR?}/samples/models/mnist.mlir"
 
 echo "=== Compiling MLIR to static library output (.vmfb, .h, .o) ==="
 ${COMPILE_TOOL?} ${INPUT_PATH} \
diff --git a/experimental/web/sample_static/iree_worker.js b/experimental/web/sample_static/iree_worker.js
index 7800669..152c34f 100644
--- a/experimental/web/sample_static/iree_worker.js
+++ b/experimental/web/sample_static/iree_worker.js
@@ -17,8 +17,10 @@
 }
 
 // TODO(scotttodd): configure this through the build system / scripts?
-const MAIN_SCRIPT_URL = 'web-sample-static-multithreaded.js';
-// const MAIN_SCRIPT_URL = 'web-sample-static-sync.js';
+// TODO(scotttodd): fix multithreading (startup silently fails on some emsdk
+//                  versions, memory growth or high initial memory also req.)
+// const MAIN_SCRIPT_URL = 'web-sample-static-multithreaded.js';
+const MAIN_SCRIPT_URL = 'web-sample-static-sync.js';
 
 let wasmSetupSampleFn;
 let wasmCleanupSampleFn;
diff --git a/experimental/web/sample_static/serve_sample.sh b/experimental/web/sample_static/serve_sample.sh
index 1688426..a9fc926 100755
--- a/experimental/web/sample_static/serve_sample.sh
+++ b/experimental/web/sample_static/serve_sample.sh
@@ -11,4 +11,4 @@
 
 echo "=== Running local webserver, open at http://localhost:8000/ ==="
 
-python3 ${ROOT_DIR?}/scripts/local_web_server.py --directory ${BINARY_DIR}
+python3 ${ROOT_DIR?}/build_tools/scripts/local_web_server.py --directory ${BINARY_DIR}
diff --git a/integrations/tensorflow/test/iree_tfl_tests/vulkan_mobilenet_v1.run b/integrations/tensorflow/test/iree_tfl_tests/vulkan_mobilenet_v1.run
index 6efa392..75129ca 100644
--- a/integrations/tensorflow/test/iree_tfl_tests/vulkan_mobilenet_v1.run
+++ b/integrations/tensorflow/test/iree_tfl_tests/vulkan_mobilenet_v1.run
@@ -1,3 +1,2 @@
 # REQUIRES: vulkan
 # RUN: %PYTHON -m iree_tfl_tests.mobilenet_v1_test --target_backend=vulkan -artifacts_dir=%t
-# XFAIL: *
diff --git a/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp b/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
index 86aef61..4bed428 100644
--- a/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
@@ -160,7 +160,8 @@
 /// Flow level.
 static SmallVector<int64_t> getDefaultDistributedLoopTileSizes(
     ArrayRef<int64_t> lbs, ArrayRef<int64_t> ubs,
-    ArrayRef<int64_t> minTileSizes, ArrayRef<int64_t> maxTileSizes) {
+    ArrayRef<int64_t> minTileSizes, ArrayRef<int64_t> maxTileSizes,
+    ArrayRef<int64_t> vectorSizeHints) {
   assert(lbs.size() == ubs.size() && lbs.size() == minTileSizes.size() &&
          lbs.size() == maxTileSizes.size() &&
          "expected all vectors to be of equal size");
@@ -181,10 +182,23 @@
     }
     int64_t candidateTileSize = 1;
     if (ubs[i] > lbs[i]) {
-      // Pick a value that evenly distributes the workload.
-      candidateTileSize = std::max<int64_t>(
-          llvm::PowerOf2Floor(static_cast<uint64_t>(ubs[i] - lbs[i]) / 2),
-          minTileSizes[i]);
+      int64_t dimSize = static_cast<uint64_t>(ubs[i] - lbs[i]);
+      int64_t targetSize = std::min(dimSize / 2, maxTileSizes[i]);
+      int64_t vectorSize = vectorSizeHints[i];
+      if (vectorSize > 1) {
+        // Pick the factor of dim which is closest to the target tile size and
+        // is a multiplier of vector size.
+        for (int64_t k = vectorSize; k <= targetSize; k += vectorSize) {
+          if (dimSize % k == 0 && k >= minTileSizes[i]) {
+            candidateTileSize = k;
+          }
+        }
+      }
+      // Fallback to power of 2 if there's no hint or can't find the ideal size.
+      if (vectorSize <= 1 || candidateTileSize == 1) {
+        candidateTileSize =
+            std::max<int64_t>(llvm::PowerOf2Floor(targetSize), minTileSizes[i]);
+      }
     }
 
     // Limit the workload per workgroup to the default being the max to keep the
@@ -205,18 +219,29 @@
   }
   unsigned currDim = numDims;
   while (numWorkgroups > numWorkgroupsLimit && currDim > 0) {
-    if (distributedTileSizes[currDim - 1] >= maxTileSizes[currDim - 1] ||
-        workload[currDim - 1] == ShapedType::kDynamicSize ||
-        distributedTileSizes[currDim - 1] >= workload[currDim - 1]) {
+    unsigned index = currDim - 1;
+    int64_t currSize = distributedTileSizes[index];
+    if (currSize >= maxTileSizes[index] ||
+        workload[index] == ShapedType::kDynamicSize ||
+        currSize >= workload[index]) {
       currDim--;
       continue;
     }
-    distributedTileSizes[currDim - 1] = std::min<int64_t>(
-        distributedTileSizes[currDim - 1] * 2, maxTileSizes[currDim - 1]);
-    int64_t nwg =
-        ceilFn(workload[currDim - 1], distributedTileSizes[currDim - 1]);
-    if (nwg < numWorkgroupsPerDim[currDim - 1]) {
-      numWorkgroups /= numWorkgroupsPerDim[currDim - 1];
+    int64_t newSize = std::min<int64_t>(currSize * 2, maxTileSizes[index]);
+    int64_t vectorSize = vectorSizeHints[index];
+    // Chech if it's the ideal size with vector size hint. And skip if the new
+    // size will break the ideal size.
+    if (vectorSize > 1 &&
+        (currSize % vectorSize == 0 && workload[index] % currSize == 0) &&
+        (newSize % vectorSize != 0 || workload[index] % newSize != 0)) {
+      currDim--;
+      continue;
+    }
+
+    distributedTileSizes[index] = newSize;
+    int64_t nwg = ceilFn(workload[index], distributedTileSizes[index]);
+    if (nwg < numWorkgroupsPerDim[index]) {
+      numWorkgroups /= numWorkgroupsPerDim[index];
       numWorkgroups *= nwg;
     } else {
       currDim--;
@@ -253,12 +278,22 @@
 
 /// Returns the tile size to use for the Flow level of an operation that
 /// implements the `PartitionableLoopsInterface`.
+///
+/// The vectorSizeHints can be empty or as many as the number of loops. When not
+/// empty, each hint should be 1 or the vector size. On the dimensions where the
+/// hints != 1, it will try to find the tile sizes which are multipliers of the
+/// hints.
 static SmallVector<int64_t> getDefaultDistributedLevelTileSizes(
     ArrayRef<Range> iterationDomain,
     IREE::Flow::PartitionableLoopsInterface partitionableLoopInterfaceOp,
-    ArrayRef<int64_t> minTileSizes, ArrayRef<int64_t> maxTileSizes) {
+    ArrayRef<int64_t> minTileSizes, ArrayRef<int64_t> maxTileSizes,
+    ArrayRef<int64_t> vectorSizeHints = {}) {
   assert(iterationDomain.size() == minTileSizes.size() &&
          "expected as many min tile sizes as number of loops");
+  assert(
+      vectorSizeHints.empty() ||
+      vectorSizeHints.size() == iterationDomain.size() &&
+          "vector size hints should be empty or equal to the number of loops");
   auto getStaticValue = [](Value v) -> int64_t {
     IntegerAttr attr;
     if (!matchPattern(v, m_Constant(&attr))) return ShapedType::kDynamicSize;
@@ -281,7 +316,8 @@
       distributedLoopUbs(numPartitionedLoops, ShapedType::kDynamicSize),
       minDistributedLoopTileSizes(numPartitionedLoops, 1),
       maxDistributedLoopTileSizes(numPartitionedLoops,
-                                  defaultWorkgroupTileSize);
+                                  defaultWorkgroupTileSize),
+      distributedVectorSizeHints(numPartitionedLoops, 1);
   // Find the bounds of the partitionable loops
   unsigned index = 0;
   for (auto range : llvm::enumerate(iterationDomain)) {
@@ -291,13 +327,15 @@
     maxDistributedLoopTileSizes[index] = maxTileSizes[range.index()];
     distributedLoopLbs[index] = lbs[range.index()];
     distributedLoopUbs[index] = ubs[range.index()];
+    distributedVectorSizeHints[index] =
+        vectorSizeHints.empty() ? 1 : vectorSizeHints[range.index()];
     index++;
   }
 
   SmallVector<int64_t> distributedTileSizes =
-      getDefaultDistributedLoopTileSizes(distributedLoopLbs, distributedLoopUbs,
-                                         minDistributedLoopTileSizes,
-                                         maxDistributedLoopTileSizes);
+      getDefaultDistributedLoopTileSizes(
+          distributedLoopLbs, distributedLoopUbs, minDistributedLoopTileSizes,
+          maxDistributedLoopTileSizes, distributedVectorSizeHints);
   SmallVector<int64_t> distributedLevelTileSizes(iterationDomain.size(), 0);
   for (auto loopID : llvm::enumerate(partitionableLoops)) {
     distributedLevelTileSizes[loopID.value()] =
@@ -827,7 +865,10 @@
   unsigned numLoops = convOp.getNumLoops();
   SmallVector<int64_t> minTileSizes(numLoops, 1);
   SmallVector<int64_t> maxTileSizes(numLoops, defaultWorkgroupTileSize);
+  SmallVector<int64_t> vectorSizeHints(numLoops, 1);
 
+  // Give the vector size hint on OC.
+  vectorSizeHints[3] = vectorSize;
   // Set the flow level tiling to the default.
   OpBuilder builder(convOp.getContext());
   builder.setInsertionPoint(convOp);
@@ -837,7 +878,7 @@
       cast<IREE::Flow::PartitionableLoopsInterface>(convOp.getOperation());
   SmallVector<int64_t> flowTileSizes = getDefaultDistributedLevelTileSizes(
       iterationDomain, partitionableLoopsInterfaceOp, minTileSizes,
-      maxTileSizes);
+      maxTileSizes, vectorSizeHints);
 
   // Shapes of N, OH, OW, OC, KH, KW, (IC)
   Optional<SmallVector<int64_t, 4>> shapes = convOp.getStaticLoopRanges();
diff --git a/iree/compiler/Codegen/LLVMCPU/Passes.cpp b/iree/compiler/Codegen/LLVMCPU/Passes.cpp
index 3748935..ee7cc49 100644
--- a/iree/compiler/Codegen/LLVMCPU/Passes.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/Passes.cpp
@@ -194,6 +194,8 @@
   passManager.addNestedPass<func::FuncOp>(createInsertDistributionInfoPass());
   passManager.addNestedPass<func::FuncOp>(
       createTileAndDistributeToWorkgroupsPass());
+  passManager.addNestedPass<func::FuncOp>(
+      createFoldAffineMinInDistributedLoopsPass());
   passManager.addPass(createCanonicalizerPass());
   passManager.addPass(createCSEPass());
 
@@ -232,6 +234,8 @@
   passManager.addNestedPass<func::FuncOp>(createInsertDistributionInfoPass());
   passManager.addNestedPass<func::FuncOp>(
       createTileAndDistributeToWorkgroupsPass());
+  passManager.addNestedPass<func::FuncOp>(
+      createFoldAffineMinInDistributedLoopsPass());
   passManager.addPass(createCanonicalizerPass());
   passManager.addPass(createCSEPass());
 
@@ -256,6 +260,8 @@
   passManager.addNestedPass<func::FuncOp>(createInsertDistributionInfoPass());
   passManager.addNestedPass<func::FuncOp>(
       createTileAndDistributeToWorkgroupsPass());
+  passManager.addNestedPass<func::FuncOp>(
+      createFoldAffineMinInDistributedLoopsPass());
   passManager.addPass(createCanonicalizerPass());
   passManager.addPass(createCSEPass());
   passManager.addNestedPass<func::FuncOp>(
@@ -324,6 +330,8 @@
   passManager.addNestedPass<func::FuncOp>(createInsertDistributionInfoPass());
   passManager.addNestedPass<func::FuncOp>(
       createTileAndDistributeToWorkgroupsPass());
+  passManager.addNestedPass<func::FuncOp>(
+      createFoldAffineMinInDistributedLoopsPass());
   passManager.addPass(createCanonicalizerPass());
   passManager.addPass(createCSEPass());
   passManager.addNestedPass<func::FuncOp>(
@@ -383,6 +391,8 @@
   passManager.addNestedPass<func::FuncOp>(createInsertDistributionInfoPass());
   passManager.addNestedPass<func::FuncOp>(
       createTileAndDistributeToWorkgroupsPass());
+  passManager.addNestedPass<func::FuncOp>(
+      createFoldAffineMinInDistributedLoopsPass());
   passManager.addPass(createCanonicalizerPass());
   passManager.addPass(createCSEPass());
 
@@ -417,6 +427,8 @@
   passManager.addNestedPass<func::FuncOp>(createInsertDistributionInfoPass());
   passManager.addNestedPass<func::FuncOp>(
       createTileAndDistributeToWorkgroupsPass());
+  passManager.addNestedPass<func::FuncOp>(
+      createFoldAffineMinInDistributedLoopsPass());
   passManager.addPass(createCanonicalizerPass());
   passManager.addPass(createCSEPass());
   // Use stack allocation on CPU side.
diff --git a/iree/compiler/Codegen/LLVMCPU/test/materialize_x86_64_launch_configuration.mlir b/iree/compiler/Codegen/LLVMCPU/test/materialize_x86_64_launch_configuration.mlir
index b8fe472..ec807f3 100644
--- a/iree/compiler/Codegen/LLVMCPU/test/materialize_x86_64_launch_configuration.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/materialize_x86_64_launch_configuration.mlir
@@ -665,7 +665,7 @@
 hal.executable private @depthwise_conv_static {
   hal.executable.variant public @system_elf_x86_64, target = <"llvm", "system-elf-x86_64", {
     data_layout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128",
-    native_vector_size = 16 : index,
+    native_vector_size = 64 : index,
     target_triple = "x86_64-unknown-linux-gnu"
   }> {
     hal.executable.entry_point public @depthwise_conv_static layout(#executable_layout)
@@ -673,27 +673,27 @@
       func.func @depthwise_conv_static() {
         %cst = arith.constant 0.0 : f32
         %input_binding = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer)
-            : !flow.dispatch.tensor<readonly:1x161x161x96xf32>
+            : !flow.dispatch.tensor<readonly:1x161x161x240xf32>
         %filter_binding = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer)
-            : !flow.dispatch.tensor<readonly:3x3x96xf32>
+            : !flow.dispatch.tensor<readonly:3x3x240xf32>
         %result_binding = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer)
-            : !flow.dispatch.tensor<writeonly:1x80x80x96xf32>
-        %input = flow.dispatch.tensor.load %input_binding, offsets = [0, 0, 0, 0], sizes = [1, 161, 161, 96], strides = [1, 1, 1, 1]
-            : !flow.dispatch.tensor<readonly:1x161x161x96xf32> -> tensor<1x161x161x96xf32>
-        %filter = flow.dispatch.tensor.load %filter_binding, offsets = [0, 0, 0], sizes = [3, 3, 96], strides = [1, 1, 1]
-            : !flow.dispatch.tensor<readonly:3x3x96xf32> -> tensor<3x3x96xf32>
-        %init = linalg.init_tensor [1, 80, 80, 96] : tensor<1x80x80x96xf32>
-        %fill = linalg.fill ins(%cst : f32) outs(%init : tensor<1x80x80x96xf32>) -> tensor<1x80x80x96xf32>
+            : !flow.dispatch.tensor<writeonly:1x80x80x240xf32>
+        %input = flow.dispatch.tensor.load %input_binding, offsets = [0, 0, 0, 0], sizes = [1, 161, 161, 240], strides = [1, 1, 1, 1]
+            : !flow.dispatch.tensor<readonly:1x161x161x240xf32> -> tensor<1x161x161x240xf32>
+        %filter = flow.dispatch.tensor.load %filter_binding, offsets = [0, 0, 0], sizes = [3, 3, 240], strides = [1, 1, 1]
+            : !flow.dispatch.tensor<readonly:3x3x240xf32> -> tensor<3x3x240xf32>
+        %init = linalg.init_tensor [1, 80, 80, 240] : tensor<1x80x80x240xf32>
+        %fill = linalg.fill ins(%cst : f32) outs(%init : tensor<1x80x80x240xf32>) -> tensor<1x80x80x240xf32>
         %conv = linalg.depthwise_conv_2d_nhwc_hwc {dilations = dense<1> : tensor<2xi64>, strides = dense<2> : tensor<2xi64>}
-            ins(%input, %filter : tensor<1x161x161x96xf32>, tensor<3x3x96xf32>) outs(%fill : tensor<1x80x80x96xf32>) -> tensor<1x80x80x96xf32>
-        flow.dispatch.tensor.store %conv, %result_binding, offsets = [0, 0, 0, 0], sizes = [1, 80, 80, 96], strides = [1, 1, 1, 1]
-            : tensor<1x80x80x96xf32> -> !flow.dispatch.tensor<writeonly:1x80x80x96xf32>
+            ins(%input, %filter : tensor<1x161x161x240xf32>, tensor<3x3x240xf32>) outs(%fill : tensor<1x80x80x240xf32>) -> tensor<1x80x80x240xf32>
+        flow.dispatch.tensor.store %conv, %result_binding, offsets = [0, 0, 0, 0], sizes = [1, 80, 80, 240], strides = [1, 1, 1, 1]
+            : tensor<1x80x80x240xf32> -> !flow.dispatch.tensor<writeonly:1x80x80x240xf32>
         return
       }
     }
   }
 }
-//  CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering_config<tile_sizes = {{\[}}[0, 20, 40, 48, 0, 0], [1, 1, 8, 8, 0, 0], [0, 0, 0, 0, 1, 3]]>
+//  CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering_config<tile_sizes = {{\[}}[0, 40, 40, 48, 0, 0], [1, 1, 8, 16, 0, 0], [0, 0, 0, 0, 1, 3]]>
 //  CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<CPUConvTileAndDecomposeExpert>
 //      CHECK: hal.executable.entry_point public @depthwise_conv_static
 // CHECK-SAME:     translation_info = #[[TRANSLATION]]
diff --git a/iree/compiler/Codegen/SPIRV/KernelConfig.cpp b/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
index 4ce42b2..a8070dd 100644
--- a/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
+++ b/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
@@ -23,6 +23,7 @@
 #include "mlir/Dialect/SPIRV/IR/TargetAndABI.h"
 #include "mlir/Dialect/Utils/StaticValueUtils.h"
 #include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/Matchers.h"
 
 #define DEBUG_TYPE "iree-spirv-kernel-config"
@@ -333,7 +334,8 @@
 //===----------------------------------------------------------------------===//
 
 static LogicalResult setDefaultOpConfig(spirv::ResourceLimitsAttr limits,
-                                        Operation *op) {
+                                        Operation *op,
+                                        bool allowVectorization = true) {
   LLVM_DEBUG(llvm::dbgs() << "Using default config for op: " << *op << "\n");
   func::FuncOp funcOp = op->getParentOfType<func::FuncOp>();
   auto interfaceOp = cast<IREE::Flow::PartitionableLoopsInterface>(*op);
@@ -407,6 +409,7 @@
   // Whether we can try to use the vectorization pipeline.
   Optional<SmallVector<int64_t, 4>> loopBounds = linalgOp.getStaticLoopRanges();
   bool vectorizable =
+      allowVectorization &&
       // The vectorization pipeline assumes tensor semantics when tiling.
       !linalgOp.hasBufferSemantics() && !linalgOp.hasIndexSemantics() &&
       // Skip vectorization for non-minor identity inputs as it generates
@@ -431,9 +434,9 @@
     // configuration for the corresponding GPU workgroup dimension.
     int64_t wgDim = 0;
     for (auto shapeDim : llvm::reverse(partitionedLoops)) {
-      // Skip untiled or dynamic dimensions.
-      // TODO: Skip size-1 dimensions in Flow level tiling and distribution.
-      if (loopBounds.getValue()[shapeDim] <= 0) continue;
+      int64_t loopBound = loopBounds.getValue()[shapeDim];
+      // Skip dynamic dimensions.
+      if (ShapedType::isDynamic(loopBound)) continue;
 
       // Try to find some power of two that can devide the current shape dim
       // size. This vector keeps the candidate tile sizes.
@@ -449,24 +452,23 @@
         candidates.push_back(i);
       }
       LLVM_DEBUG({
-        llvm::dbgs() << "Candidates tile sizes: [";
+        llvm::dbgs() << "Candidate tile sizes: [";
         llvm::interleaveComma(candidates, llvm::dbgs());
         llvm::dbgs() << "]\n";
       });
 
       for (int64_t candidate : candidates) {
-        if (loopBounds.getValue()[shapeDim] % candidate != 0) {
+        if (loopBound % candidate != 0) {
           if (!lossFactor) continue;
           // Skip this candidate if it causes many threads to be idle.
-          int64_t idleThreads =
-              candidate - (loopBounds.getValue()[shapeDim] % candidate);
+          int64_t idleThreads = candidate - (loopBound % candidate);
           if (idleThreads > candidate / *lossFactor) continue;
         }
-        LLVM_DEBUG(llvm::dbgs() << "Chosen Candiate " << candidate << "\n");
 
         // Found a suitable candidate. Try to let each thread handle 4
         // elements if this is the workgroup x dimension.
         workgroupTileSizes[shapeDim] = candidate;
+        LLVM_DEBUG(llvm::dbgs() << "Chosen tile size: " << candidate << "\n");
         if (vectorizable && wgDim == 0 && !lossFactor && candidate % 4 == 0) {
           threadTileSizes[shapeDim] = 4;
           workgroupSize[wgDim] = candidate / 4;
@@ -578,8 +580,9 @@
             // If unsuccessful, try to tile and distribute.
             return setDefaultOpConfig(limits, op);
           })
-      .Case<IREE::LinalgExt::FftOp>([limits](IREE::LinalgExt::FftOp op) {
-        return setFftOpConfig(limits, op);
+      .Case<linalg::ConvolutionOpInterface>([limits](auto op) {
+        // Other convolution/pooling op vectorization is not wired up.
+        return setDefaultOpConfig(limits, op, /*allowVectorization=*/false);
       })
       .Case<linalg::GenericOp>([limits](linalg::GenericOp op) {
         // If a generic op has reduction iterator types, it can be treated as a
@@ -590,6 +593,9 @@
         }
         return success();
       })
+      .Case<IREE::LinalgExt::FftOp>([limits](IREE::LinalgExt::FftOp op) {
+        return setFftOpConfig(limits, op);
+      })
       .Default([](Operation *) { return success(); });
 };
 
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVVectorize.cpp b/iree/compiler/Codegen/SPIRV/SPIRVVectorize.cpp
index fd2ae6b..3808661 100644
--- a/iree/compiler/Codegen/SPIRV/SPIRVVectorize.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVVectorize.cpp
@@ -100,7 +100,6 @@
   patterns.add<linalg::LinalgVectorizationPattern>(
       patterns.getContext(), f.addOpFilter<linalg::ContractionOpInterface>(),
       opt);
-  populateVectorizePadPatterns(patterns);
   vector::populateVectorTransferPermutationMapLoweringPatterns(patterns);
   vector::populateVectorReductionToContractPatterns(patterns);
 }
@@ -131,7 +130,9 @@
     {
       RewritePatternSet patterns(context);
       populateVectorizationPatterns(patterns);
+      // Pull in additional vectorization patterns in IREE.
       populateLinalgToVectorVectorizeConvPatterns(context, patterns);
+      populateVectorizePadPatterns(patterns);
       if (failed(applyPatternsAndFoldGreedily(funcOp, std::move(patterns)))) {
         return signalPassFailure();
       }
diff --git a/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ext_ops.mlir b/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ext_ops.mlir
index 7dc3add..e8727a5 100644
--- a/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ext_ops.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ext_ops.mlir
@@ -184,47 +184,3 @@
 //       CHECK: func @static_3d_fft_stage3()
 //       CHECK:   iree_linalg_ext.fft
 //  CHECK-SAME:     lowering_config = #[[CONFIG]]
-
-// -----
-
-#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
-  #hal.descriptor_set.layout<0, bindings = [
-    #hal.descriptor_set.binding<0, storage_buffer>,
-    #hal.descriptor_set.binding<1, storage_buffer>
-  ]>
-]>
-hal.executable private @copy_op {
-  hal.executable.variant @vulkan_spirv_fb, target = <"vulkan", "vulkan-spirvfb", {
-      spv.target_env = #spv.target_env<#spv.vce<v1.4, [Shader], []>, Unknown:IntegratedGPU, {
-        max_compute_shared_memory_size = 32768 : i32,
-        max_compute_workgroup_invocations = 512 : i32,
-        max_compute_workgroup_size = dense<512> : vector<3xi32>,
-        subgroup_size = 16 : i32}>
-    }> {
-    hal.executable.entry_point @copy_op layout(#executable_layout)
-    builtin.module {
-      func.func @copy_op() {
-        %offset_y = hal.interface.constant.load[0] : index
-        %offset_x = hal.interface.constant.load[1] : index
-        %source_size_y = hal.interface.constant.load[2] : index
-        %source_size_x = hal.interface.constant.load[3] : index
-        %dest_size_y = hal.interface.constant.load[4] : index
-        %dest_size_x = hal.interface.constant.load[5] : index
-        %source = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : memref<?x?xf32>{%source_size_y, %source_size_x}
-        %dest = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : memref<?x?xf32>{%dest_size_y, %dest_size_x}
-        linalg.generic {
-            indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>],
-            iterator_types = ["parallel", "parallel"]}
-            ins(%source : memref<?x?xf32>) outs(%dest : memref<?x?xf32>) {
-          ^bb0(%b0 : f32, %b1 : f32):
-            linalg.yield %b0 : f32
-          }
-        return
-      }
-    }
-  }
-}
-//  CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering_config<tile_sizes = {{\[}}[1, 16], [1, 1]{{\]}}>
-//  CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<SPIRVDistribute>
-//      CHECK: linalg.generic
-// CHECK-SAME:     lowering_config = #[[CONFIG]]
diff --git a/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ops.mlir b/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ops.mlir
index 034ff39..a25cd54 100644
--- a/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ops.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ops.mlir
@@ -118,7 +118,7 @@
             : !flow.dispatch.tensor<readonly:1x24x24x8xf32> -> tensor<1x24x24x8xf32>
         %20 = linalg.init_tensor [1, 2, 2, 8] : tensor<1x2x2x8xf32>
         %21 = linalg.fill ins(%cst : f32) outs(%20 : tensor<1x2x2x8xf32>) -> tensor<1x2x2x8xf32>
-        %22 = linalg.pooling_nhwc_sum {__internal_linalg_transform__ = "workgroup", dilations = dense<1> : vector<2xi64>, strides = dense<12> : vector<2xi64>}
+        %22 = linalg.pooling_nhwc_sum {dilations = dense<1> : vector<2xi64>, strides = dense<12> : vector<2xi64>}
             ins(%14, %2 : tensor<1x24x24x8xf32>, tensor<12x12xf32>)
             outs(%21 : tensor<1x2x2x8xf32>) -> tensor<1x2x2x8xf32>
         flow.dispatch.tensor.store %22, %1, offsets = [0, 0, 0, 0], sizes = [1, 2, 2, 8], strides = [1, 1, 1, 1]
@@ -138,6 +138,63 @@
 
 // -----
 
+// Polling vectorization is not supported for now.
+
+#executable_layout = #hal.executable.layout<push_constants = 0, sets = [
+  #hal.descriptor_set.layout<0, bindings = [
+    #hal.descriptor_set.binding<0, storage_buffer>,
+    #hal.descriptor_set.binding<1, storage_buffer>
+  ]>
+]>
+hal.executable @avg_pool {
+  hal.executable.variant @vulkan_spirv_fb, target = <"vulkan-spirv", "vulkan-spirv-fb", {
+      spv.target_env = #spv.target_env<#spv.vce<v1.4, [Shader], []>, Unknown:IntegratedGPU, {
+        max_compute_shared_memory_size = 32768 : i32,
+        max_compute_workgroup_invocations = 512 : i32,
+        max_compute_workgroup_size = dense<512> : vector<3xi32>,
+        subgroup_size = 4 : i32}>
+    }> {
+    hal.executable.entry_point public @avg_pool layout(#executable_layout)
+    builtin.module {
+      func.func @avg_pool() {
+        %cst = arith.constant 0.000000e+00 : f32
+        %cst_0 = arith.constant 4.900000e+01 : f32
+        %c0 = arith.constant 0 : index
+        %0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) offset(%c0) alignment(64) : !flow.dispatch.tensor<readonly:1x7x7x1280xf32>
+        %1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) offset(%c0) alignment(64) : !flow.dispatch.tensor<writeonly:1x1x1x1280xf32>
+        %2 = flow.dispatch.tensor.load %0, offsets = [0, 0, 0, 0], sizes = [1, 7, 7, 1280], strides = [1, 1, 1, 1]
+          : !flow.dispatch.tensor<readonly:1x7x7x1280xf32> -> tensor<1x7x7x1280xf32>
+        %3 = linalg.init_tensor [7, 7] : tensor<7x7xf32>
+        %4 = linalg.init_tensor [1, 1, 1, 1280] : tensor<1x1x1x1280xf32>
+        %5 = linalg.fill ins(%cst : f32) outs(%4 : tensor<1x1x1x1280xf32>) -> tensor<1x1x1x1280xf32>
+        %6 = linalg.pooling_nhwc_sum {
+          dilations = dense<1> : vector<2xi64>, strides = dense<1> : vector<2xi64>
+        } ins(%2, %3 : tensor<1x7x7x1280xf32>, tensor<7x7xf32>) outs(%5 : tensor<1x1x1x1280xf32>) -> tensor<1x1x1x1280xf32>
+        %7 = linalg.generic {
+          indexing_maps = [affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>, affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>],
+          iterator_types = ["parallel", "parallel", "parallel", "parallel"]
+        } ins(%6 : tensor<1x1x1x1280xf32>) outs(%4 : tensor<1x1x1x1280xf32>) {
+        ^bb0(%arg0: f32, %arg1: f32):
+          %8 = arith.divf %arg0, %cst_0 : f32
+          linalg.yield %8 : f32
+        } -> tensor<1x1x1x1280xf32>
+        flow.dispatch.tensor.store %7, %1, offsets = [0, 0, 0, 0], sizes = [1, 1, 1, 1280], strides = [1, 1, 1, 1]
+          : tensor<1x1x1x1280xf32> -> !flow.dispatch.tensor<writeonly:1x1x1x1280xf32>
+        return
+      }
+    }
+  }
+}
+
+//  CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering_config<tile_sizes = {{\[}}[0, 0, 0, 4], [0, 0, 0, 1]{{\]}}>
+//  CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<SPIRVDistribute>
+//      CHECK: hal.executable.entry_point public @avg_pool
+// CHECK-SAME:   translation_info = #[[TRANSLATION]]
+//      CHECK:   linalg.pooling_nhwc_sum
+// CHECK-SAME:     lowering_config = #[[CONFIG]]
+
+// -----
+
 // Max pooling op with odd size-1 dimension sizes.
 
 #executable_layout = #hal.executable.layout<push_constants = 0, sets = [
diff --git a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/BUILD b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/BUILD
index affc8b3..70d1c73 100644
--- a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/BUILD
+++ b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/BUILD
@@ -42,6 +42,7 @@
             "conversion_ops_f32.mlir",
             "conversion_ops_i64.mlir",
             "conversion_ops.mlir",
+            "func_op.mlir",
             "global_ops_f32.mlir",
             "global_ops_i64.mlir",
             "global_ops.mlir",
diff --git a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/CMakeLists.txt b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/CMakeLists.txt
index 01c8c56..63d7152 100644
--- a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/CMakeLists.txt
+++ b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/CMakeLists.txt
@@ -34,6 +34,7 @@
     "conversion_ops.mlir"
     "conversion_ops_f32.mlir"
     "conversion_ops_i64.mlir"
+    "func_op.mlir"
     "global_ops.mlir"
     "global_ops_f32.mlir"
     "global_ops_i64.mlir"
diff --git a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/control_flow_ops.mlir b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/control_flow_ops.mlir
index d4f939a..c12a755 100644
--- a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/control_flow_ops.mlir
+++ b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/control_flow_ops.mlir
@@ -234,6 +234,210 @@
 
 // -----
 
+// Test vm.import conversion on a variadic function.
+vm.module @my_module {
+
+  // CHECK-LABEL: func @my_module_call_0iCiD_i_2_import_shim(%arg0: !emitc.ptr<!emitc.opaque<"iree_vm_stack_t">>, %arg1: !emitc.ptr<!emitc.opaque<"iree_vm_function_t">>,
+  // CHECK-SAME:                                             %arg2: i32, %arg3: i32, %arg4: i32, %arg5: i32, %arg6: !emitc.ptr<!emitc.opaque<"int32_t">>)
+  // CHECK-SAME:      -> !emitc.opaque<"iree_status_t"> attributes {emitc.static} {
+
+  // Calculate the size of the arguments.
+  // CHECK-NEXT: %[[ARGSIZE0:.+]] = "emitc.constant"() {value = #emitc.opaque<"0">} : () -> !emitc.opaque<"iree_host_size_t">
+  // CHECK-NEXT: %[[ARGSIZE1:.+]] = emitc.call "sizeof"() {args = [#emitc.opaque<"int32_t">]}
+  // CHECK-NEXT: %[[ARGSIZE01:.+]] = emitc.call "EMITC_ADD"(%[[ARGSIZE0]], %[[ARGSIZE1]])
+  // CHECK-NEXT: %[[ARGSIZE2:.+]] = emitc.call "sizeof"() {args = [#emitc.opaque<"int32_t">]}
+  // CHECK-NEXT: %[[ARGSIZE012:.+]] = emitc.call "EMITC_ADD"(%[[ARGSIZE01]], %[[ARGSIZE2]])
+  // CHECK-NEXT: %[[ARGSIZE3:.+]] = emitc.call "sizeof"() {args = [#emitc.opaque<"int32_t">]}
+  // CHECK-NEXT: %[[ARGSIZE0123:.+]] = emitc.call "EMITC_ADD"(%[[ARGSIZE012]], %[[ARGSIZE3]])
+  // CHECK-NEXT: %[[ARGSIZE4:.+]] = emitc.call "sizeof"() {args = [#emitc.opaque<"int32_t">]}
+  // CHECK-NEXT: %[[ARGSIZE:.+]] = emitc.call "EMITC_ADD"(%[[ARGSIZE0123]], %[[ARGSIZE4]])
+
+  // Calculate the size of the result.
+  // CHECK-NEXT: %[[RESULTSIZE0:.+]] = "emitc.constant"() {value = #emitc.opaque<"0">} : () -> !emitc.opaque<"iree_host_size_t">
+  // CHECK-NEXT: %[[RESULTSIZE1:.+]] = emitc.call "sizeof"() {args = [#emitc.opaque<"int32_t">]}
+  // CHECK-NEXT: %[[RESULTSIZE:.+]] = emitc.call "EMITC_ADD"(%[[RESULTSIZE0]], %[[RESULTSIZE1]])
+
+  // Create a struct for the arguments and results.
+  // CHECK: %[[ARGSTRUCT:.+]] = "emitc.constant"() {value = #emitc.opaque<"">} : () -> !emitc.opaque<"iree_vm_function_call_t">
+  // CHECK-NEXT: %[[ARGSTRUCTFN:.+]] = emitc.apply "*"(%arg1) : (!emitc.ptr<!emitc.opaque<"iree_vm_function_t">>) -> !emitc.opaque<"iree_vm_function_t">
+  // CHECK-NEXT: emitc.call "EMITC_STRUCT_MEMBER_ASSIGN"(%[[ARGSTRUCT]], %[[ARGSTRUCTFN]]) {args = [0 : index, #emitc.opaque<"function">, 1 : index]}
+
+  // Allocate space for the arguments.
+  // CHECK-NEXT: %[[ARGBYTESPAN:.+]] = emitc.call "EMITC_STRUCT_MEMBER_ADDRESS"(%[[ARGSTRUCT]]) {args = [0 : index, #emitc.opaque<"arguments">]}
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_function_call_t">) -> !emitc.ptr<!emitc.opaque<"iree_byte_span_t">>
+  // CHECK-NEXT: %[[ARGBYTESPANDATAVOID:.+]] = emitc.call "iree_alloca"(%[[ARGSIZE]])
+  // CHECK-SAME:     : (!emitc.opaque<"iree_host_size_t">) -> !emitc.ptr<!emitc.opaque<"void">>
+  // CHECK-NEXT: %[[ARGBYTESPANDATA:.+]] = emitc.call "EMITC_CAST"(%[[ARGBYTESPANDATAVOID]]) {args = [0 : index, !emitc.ptr<ui8>]}
+  // CHECK-SAME:     : (!emitc.ptr<!emitc.opaque<"void">>) -> !emitc.ptr<ui8>
+  // CHECK-NEXT: emitc.call "EMITC_STRUCT_PTR_MEMBER_ASSIGN"(%[[ARGBYTESPAN]], %[[ARGSIZE]]) {args = [0 : index, #emitc.opaque<"data_length">, 1 : index]}
+  // CHECK-NEXT: emitc.call "EMITC_STRUCT_PTR_MEMBER_ASSIGN"(%[[ARGBYTESPAN]], %[[ARGBYTESPANDATA]]) {args = [0 : index, #emitc.opaque<"data">, 1 : index]}
+  // CHECK-NEXT: emitc.call "memset"(%[[ARGBYTESPANDATA]], %[[ARGSIZE]]) {args = [0 : index, 0 : i32, 1 : index]}
+
+  // Allocate space for the result.
+  // CHECK-NEXT: %[[RESBYTESPAN:.+]] = emitc.call "EMITC_STRUCT_MEMBER_ADDRESS"(%[[ARGSTRUCT]]) {args = [0 : index, #emitc.opaque<"results">]}
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_function_call_t">) -> !emitc.ptr<!emitc.opaque<"iree_byte_span_t">>
+  // CHECK-NEXT: %[[RESBYTESPANDATAVOID:.+]] = emitc.call "iree_alloca"(%[[RESULTSIZE]])
+  // CHECK-SAME:     : (!emitc.opaque<"iree_host_size_t">) -> !emitc.ptr<!emitc.opaque<"void">>
+  // CHECK-NEXT: %[[RESBYTESPANDATA:.+]] = emitc.call "EMITC_CAST"(%[[RESBYTESPANDATAVOID]]) {args = [0 : index, !emitc.ptr<ui8>]}
+  // CHECK-SAME:     : (!emitc.ptr<!emitc.opaque<"void">>) -> !emitc.ptr<ui8>
+  // CHECK-NEXT: emitc.call "EMITC_STRUCT_PTR_MEMBER_ASSIGN"(%[[RESBYTESPAN]], %[[RESULTSIZE]]) {args = [0 : index, #emitc.opaque<"data_length">, 1 : index]}
+  // CHECK-NEXT: emitc.call "EMITC_STRUCT_PTR_MEMBER_ASSIGN"(%[[RESBYTESPAN]], %[[RESBYTESPANDATA]]) {args = [0 : index, #emitc.opaque<"data">, 1 : index]}
+  // CHECK-NEXT: emitc.call "memset"(%[[RESBYTESPANDATA]], %[[RESULTSIZE]]) {args = [0 : index, 0 : i32, 1 : index]}
+
+  // Pack the arguments into the struct.
+  // Here we also create pointers for non-pointer types.
+  // CHECK-NEXT: %[[ARGS:.+]] = emitc.call "EMITC_STRUCT_MEMBER"(%[[ARGSTRUCT]]) {args = [0 : index, #emitc.opaque<"arguments">]}
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_function_call_t">) -> !emitc.opaque<"iree_byte_span_t">
+  // CHECK-NEXT: %[[ARGSPTR:.+]] = emitc.call "EMITC_STRUCT_MEMBER"(%[[ARGS]]) {args = [0 : index, #emitc.opaque<"data">]}
+  // CHECK-SAME:     : (!emitc.opaque<"iree_byte_span_t">) -> !emitc.ptr<ui8>
+  // CHECK-NEXT: %[[ARGHOSTSIZE:.+]] = emitc.call "sizeof"() {args = [#emitc.opaque<"int32_t">]} : () -> !emitc.opaque<"iree_host_size_t">
+  // CHECK-NEXT: %[[A1PTR:.+]] = emitc.apply "&"(%arg2) : (i32) -> !emitc.ptr<!emitc.opaque<"int32_t">>
+  // CHECK-NEXT: emitc.call "memcpy"(%[[ARGSPTR]], %[[A1PTR]], %[[ARGHOSTSIZE]])
+  // CHECK-NEXT: %[[A1ADR:.+]] = emitc.call "EMITC_ADD"(%[[ARGSPTR]], %[[ARGHOSTSIZE]])
+  // CHECK-SAME:     : (!emitc.ptr<ui8>, !emitc.opaque<"iree_host_size_t">) -> !emitc.ptr<ui8>
+  // CHECK-NEXT: %[[A1SIZE:.+]] = emitc.call "sizeof"() {args = [#emitc.opaque<"int32_t">]} : () -> !emitc.opaque<"iree_host_size_t">
+  // CHECK-NEXT: %[[A2PTR:.+]] = emitc.apply "&"(%arg3) : (i32) -> !emitc.ptr<!emitc.opaque<"int32_t">>
+  // CHECK-NEXT: emitc.call "memcpy"(%[[A1ADR]], %[[A2PTR]], %[[A1SIZE]])
+  // CHECK-NEXT: %[[A2ADR:.+]] = emitc.call "EMITC_ADD"(%[[A1ADR]], %[[A1SIZE]])
+  // CHECK-SAME:     : (!emitc.ptr<ui8>, !emitc.opaque<"iree_host_size_t">) -> !emitc.ptr<ui8>
+  // CHECK-NEXT: %[[A2SIZE:.+]] = emitc.call "sizeof"() {args = [#emitc.opaque<"int32_t">]} : () -> !emitc.opaque<"iree_host_size_t">
+  // CHECK-NEXT: %[[A3PTR:.+]] = emitc.apply "&"(%arg4) : (i32) -> !emitc.ptr<!emitc.opaque<"int32_t">>
+  // CHECK-NEXT: emitc.call "memcpy"(%[[A2ADR]], %[[A3PTR]], %[[A2SIZE]])
+  // CHECK-NEXT: %[[A3ADR:.+]] = emitc.call "EMITC_ADD"(%[[A2ADR]], %[[A2SIZE]])
+  // CHECK-SAME:     : (!emitc.ptr<ui8>, !emitc.opaque<"iree_host_size_t">) -> !emitc.ptr<ui8>
+  // CHECK-NEXT: %[[A3SIZE:.+]] = emitc.call "sizeof"() {args = [#emitc.opaque<"int32_t">]} : () -> !emitc.opaque<"iree_host_size_t">
+  // CHECK-NEXT: %[[A4PTR:.+]] = emitc.apply "&"(%arg5) : (i32) -> !emitc.ptr<!emitc.opaque<"int32_t">>
+  // CHECK-NEXT: emitc.call "memcpy"(%[[A3ADR]], %[[A4PTR]], %[[A3SIZE:.+]])
+
+  // Create the call to the imported function.
+  // CHECK-NEXT: %[[EXECRESULT:.+]] = "emitc.variable"() {value = #emitc.opaque<"">} : () -> !emitc.opaque<"iree_vm_execution_result_t">
+  // CHECK-NEXT: %[[EXECRESULTPTR1:.+]] = emitc.apply "&"(%[[EXECRESULT]])
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_execution_result_t">) -> !emitc.ptr<!emitc.opaque<"iree_vm_execution_result_t">>
+  // CHECK-NEXT: %[[EXECRESULTSIZE:.+]] = emitc.call "sizeof"(%[[EXECRESULT]]) : (!emitc.opaque<"iree_vm_execution_result_t">) -> i32
+  // CHECK-NEXT: emitc.call "memset"(%[[EXECRESULTPTR1]], %[[EXECRESULTSIZE]]) {args = [0 : index, 0 : ui32, 1 : index]}
+  // CHECK-SAME:     : (!emitc.ptr<!emitc.opaque<"iree_vm_execution_result_t">>, i32) -> ()
+  // CHECK-NEXT: %[[IMPORTMOD:.+]] = emitc.call "EMITC_STRUCT_PTR_MEMBER"(%arg1) {args = [0 : index, #emitc.opaque<"module">]}
+  // CHECK-SAME:     : (!emitc.ptr<!emitc.opaque<"iree_vm_function_t">>) -> !emitc.ptr<!emitc.opaque<"iree_vm_module_t">>
+  // CHECK-NEXT: %[[ARGSTRUCTPTR:.+]] = emitc.apply "&"(%[[ARGSTRUCT]])
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_function_call_t">) -> !emitc.ptr<!emitc.opaque<"iree_vm_function_call_t">>
+  // CHECK-NEXT: %[[EXECRESULTPTR:.+]] = emitc.apply "&"(%[[EXECRESULT]])
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_execution_result_t">) -> !emitc.ptr<!emitc.opaque<"iree_vm_execution_result_t">>
+  // CHECK-NEXT: %{{.+}} = emitc.call "EMITC_STRUCT_PTR_MEMBER_CALL"(%[[IMPORTMOD]], %arg0, %[[ARGSTRUCTPTR]], %[[EXECRESULTPTR]])
+  // CHECK-SAME:     {args = [0 : index, #emitc.opaque<"begin_call">, 0 : index, 1 : index, 2 : index, 3 : index]}
+
+  // Unpack the function results.
+  //      CHECK: %[[RES:.+]] = emitc.call "EMITC_STRUCT_MEMBER"(%[[ARGSTRUCT]]) {args = [0 : index, #emitc.opaque<"results">]}
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_function_call_t">) -> !emitc.opaque<"iree_byte_span_t">
+  // CHECK-NEXT: %[[RESPTR:.+]] = emitc.call "EMITC_STRUCT_MEMBER"(%[[RES]]) {args = [0 : index, #emitc.opaque<"data">]}
+  // CHECK-SAME:     : (!emitc.opaque<"iree_byte_span_t">) -> !emitc.ptr<ui8>
+  // CHECK-NEXT: %[[RESHOSTSIZE:.+]] = emitc.call "sizeof"() {args = [#emitc.opaque<"int32_t">]} : () -> !emitc.opaque<"iree_host_size_t">
+  // CHECK-NEXT: emitc.call "memcpy"(%arg6, %[[RESPTR]], %[[RESHOSTSIZE]])
+
+  // Return ok status.
+  // CHECK-NEXT: %[[OK:.+]] = emitc.call "iree_ok_status"()
+  // CHECK-NEXT: return %[[OK]]
+  vm.import @variadic_fn(%arg0 : i32, %arg1 : i32 ...) -> i32
+
+  vm.func @import_variadic(%arg0 : i32, %arg1 : i32, %arg2 : i32) -> i32 {
+    %0 = vm.call.variadic @variadic_fn(%arg0, [%arg1, %arg2]) : (i32, i32 ...) -> i32
+    vm.return %0 : i32
+  }
+}
+
+// -----
+
+// Test vm.import conversion on a function with vm.ref arguments.
+vm.module @my_module {
+
+  // CHECK-LABEL: func @my_module_call_0r_r_import_shim(%arg0: !emitc.ptr<!emitc.opaque<"iree_vm_stack_t">>, %arg1: !emitc.ptr<!emitc.opaque<"iree_vm_function_t">>,
+  // CHECK-SAME:                                        %arg2: !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, %arg3: !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>)
+  // CHECK-SAME:      -> !emitc.opaque<"iree_status_t"> attributes {emitc.static} {
+
+  // Calculate the size of the arguments.
+  // CHECK-NEXT: %[[ARGSIZE0:.+]] = "emitc.constant"() {value = #emitc.opaque<"0">} : () -> !emitc.opaque<"iree_host_size_t">
+  // CHECK-NEXT: %[[ARGSIZE1:.+]] = emitc.call "sizeof"() {args = [#emitc.opaque<"iree_vm_ref_t">]}
+  // CHECK-NEXT: %[[ARGSIZE:.+]] = emitc.call "EMITC_ADD"(%[[ARGSIZE0]], %[[ARGSIZE1]])
+
+  // Calculate the size of the result.
+  // CHECK-NEXT: %[[RESULTSIZE0:.+]] = "emitc.constant"() {value = #emitc.opaque<"0">} : () -> !emitc.opaque<"iree_host_size_t">
+  // CHECK-NEXT: %[[RESULTSIZE1:.+]] = emitc.call "sizeof"() {args = [#emitc.opaque<"iree_vm_ref_t">]}
+  // CHECK-NEXT: %[[RESULTSIZE:.+]] = emitc.call "EMITC_ADD"(%[[RESULTSIZE0]], %[[RESULTSIZE1]])
+
+  // Create a struct for the arguments and results.
+  // CHECK: %[[ARGSTRUCT:.+]] = "emitc.constant"() {value = #emitc.opaque<"">} : () -> !emitc.opaque<"iree_vm_function_call_t">
+  // CHECK-NEXT: %[[ARGSTRUCTFN:.+]] = emitc.apply "*"(%arg1) : (!emitc.ptr<!emitc.opaque<"iree_vm_function_t">>) -> !emitc.opaque<"iree_vm_function_t">
+  // CHECK-NEXT: emitc.call "EMITC_STRUCT_MEMBER_ASSIGN"(%[[ARGSTRUCT]], %[[ARGSTRUCTFN]]) {args = [0 : index, #emitc.opaque<"function">, 1 : index]}
+
+  // Allocate space for the arguments.
+  // CHECK-NEXT: %[[ARGBYTESPAN:.+]] = emitc.call "EMITC_STRUCT_MEMBER_ADDRESS"(%[[ARGSTRUCT]]) {args = [0 : index, #emitc.opaque<"arguments">]}
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_function_call_t">) -> !emitc.ptr<!emitc.opaque<"iree_byte_span_t">>
+  // CHECK-NEXT: %[[ARGBYTESPANDATAVOID:.+]] = emitc.call "iree_alloca"(%[[ARGSIZE]]) : (!emitc.opaque<"iree_host_size_t">) -> !emitc.ptr<!emitc.opaque<"void">>
+  // CHECK-NEXT: %[[ARGBYTESPANDATA:.+]] = emitc.call "EMITC_CAST"(%[[ARGBYTESPANDATAVOID]]) {args = [0 : index, !emitc.ptr<ui8>]}
+  // CHECK-SAME:     : (!emitc.ptr<!emitc.opaque<"void">>) -> !emitc.ptr<ui8>
+  // CHECK-NEXT: emitc.call "EMITC_STRUCT_PTR_MEMBER_ASSIGN"(%[[ARGBYTESPAN]], %[[ARGSIZE]]) {args = [0 : index, #emitc.opaque<"data_length">, 1 : index]}
+  // CHECK-NEXT: emitc.call "EMITC_STRUCT_PTR_MEMBER_ASSIGN"(%[[ARGBYTESPAN]], %[[ARGBYTESPANDATA]]) {args = [0 : index, #emitc.opaque<"data">, 1 : index]}
+  // CHECK-NEXT: emitc.call "memset"(%[[ARGBYTESPANDATA]], %[[ARGSIZE]]) {args = [0 : index, 0 : i32, 1 : index]}
+
+  // Allocate space for the result.
+  // CHECK-NEXT: %[[RESBYTESPAN:.+]] = emitc.call "EMITC_STRUCT_MEMBER_ADDRESS"(%[[ARGSTRUCT]]) {args = [0 : index, #emitc.opaque<"results">]}
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_function_call_t">) -> !emitc.ptr<!emitc.opaque<"iree_byte_span_t">>
+  // CHECK-NEXT: %[[RESBYTESPANDATAVOID:.+]] = emitc.call "iree_alloca"(%[[RESULTSIZE]]) : (!emitc.opaque<"iree_host_size_t">) -> !emitc.ptr<!emitc.opaque<"void">>
+  // CHECK-NEXT: %[[RESBYTESPANDATA:.+]] = emitc.call "EMITC_CAST"(%[[RESBYTESPANDATAVOID]]) {args = [0 : index, !emitc.ptr<ui8>]}
+  // CHECK-SAME:     : (!emitc.ptr<!emitc.opaque<"void">>) -> !emitc.ptr<ui8>
+  // CHECK-NEXT: emitc.call "EMITC_STRUCT_PTR_MEMBER_ASSIGN"(%[[RESBYTESPAN]], %[[RESULTSIZE]]) {args = [0 : index, #emitc.opaque<"data_length">, 1 : index]}
+  // CHECK-NEXT: emitc.call "EMITC_STRUCT_PTR_MEMBER_ASSIGN"(%[[RESBYTESPAN]], %[[RESBYTESPANDATA]]) {args = [0 : index, #emitc.opaque<"data">, 1 : index]}
+  // CHECK-NEXT: emitc.call "memset"(%[[RESBYTESPANDATA]], %[[RESULTSIZE]]) {args = [0 : index, 0 : i32, 1 : index]}
+
+  // Pack the argument into the struct.
+  // CHECK-NEXT: %[[ARGS:.+]] = emitc.call "EMITC_STRUCT_MEMBER"(%[[ARGSTRUCT]]) {args = [0 : index, #emitc.opaque<"arguments">]}
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_function_call_t">) -> !emitc.opaque<"iree_byte_span_t">
+  // CHECK-NEXT: %[[ARGSPTR:.+]] = emitc.call "EMITC_STRUCT_MEMBER"(%[[ARGS]]) {args = [0 : index, #emitc.opaque<"data">]}
+  // CHECK-SAME:     : (!emitc.opaque<"iree_byte_span_t">) -> !emitc.ptr<ui8>
+  // CHECK-NEXT: %{{.+}} = emitc.call "sizeof"() {args = [#emitc.opaque<"iree_vm_ref_t">]} : () -> !emitc.opaque<"iree_host_size_t">
+  // CHECK-NEXT: %[[ARG:.+]] = emitc.call "EMITC_CAST"(%[[ARGSPTR]]) {args = [0 : index, !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>]}
+  // CHECK-SAME:     : (!emitc.ptr<ui8>) -> !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>
+  // CHECK-NEXT: emitc.call "iree_vm_ref_assign"(%arg2, %[[ARG]])
+
+  // Create the call to the imported function.
+  // CHECK-NEXT: %[[EXECRESULT:.+]] = "emitc.variable"() {value = #emitc.opaque<"">}
+  // CHECK-SAME:     : () -> !emitc.opaque<"iree_vm_execution_result_t">
+  // CHECK-NEXT: %[[EXECRESULTPTR1:.+]] = emitc.apply "&"(%[[EXECRESULT]])
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_execution_result_t">) -> !emitc.ptr<!emitc.opaque<"iree_vm_execution_result_t">>
+  // CHECK-NEXT: %[[EXECRESULTSIZE:.+]] = emitc.call "sizeof"(%[[EXECRESULT]])
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_execution_result_t">) -> i32
+  // CHECK-NEXT: emitc.call "memset"(%[[EXECRESULTPTR1]], %[[EXECRESULTSIZE]]) {args = [0 : index, 0 : ui32, 1 : index]}
+  // CHECK-NEXT: %[[IMPORTMOD:.+]] = emitc.call "EMITC_STRUCT_PTR_MEMBER"(%arg1) {args = [0 : index, #emitc.opaque<"module">]}
+  // CHECK-SAME:     : (!emitc.ptr<!emitc.opaque<"iree_vm_function_t">>) -> !emitc.ptr<!emitc.opaque<"iree_vm_module_t">>
+  // CHECK-NEXT: %[[ARGSTRUCTPTR:.+]] = emitc.apply "&"(%[[ARGSTRUCT]])
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_function_call_t">) -> !emitc.ptr<!emitc.opaque<"iree_vm_function_call_t">>
+  // CHECK-NEXT: %[[EXECRESULTPTR:.+]] = emitc.apply "&"(%[[EXECRESULT]])
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_execution_result_t">) -> !emitc.ptr<!emitc.opaque<"iree_vm_execution_result_t">>
+  // CHECK-NEXT: %{{.+}} = emitc.call "EMITC_STRUCT_PTR_MEMBER_CALL"(%[[IMPORTMOD]], %arg0, %[[ARGSTRUCTPTR]], %[[EXECRESULTPTR]])
+  // CHECK-SAME:     {args = [0 : index, #emitc.opaque<"begin_call">, 0 : index, 1 : index, 2 : index, 3 : index]}
+
+  // Unpack the function results.
+  //      CHECK: %[[RES:.+]] = emitc.call "EMITC_STRUCT_MEMBER"(%[[ARGSTRUCT]]) {args = [0 : index, #emitc.opaque<"results">]}
+  // CHECK-SAME:     : (!emitc.opaque<"iree_vm_function_call_t">) -> !emitc.opaque<"iree_byte_span_t">
+  // CHECK-NEXT: %[[RESPTR:.+]] = emitc.call "EMITC_STRUCT_MEMBER"(%[[RES]]) {args = [0 : index, #emitc.opaque<"data">]}
+  // CHECK-SAME:     : (!emitc.opaque<"iree_byte_span_t">) -> !emitc.ptr<ui8>
+  // CHECK-NEXT: %{{.+}} = emitc.call "sizeof"() {args = [#emitc.opaque<"iree_vm_ref_t">]} : () -> !emitc.opaque<"iree_host_size_t">
+  // CHECK-NEXT: %[[RESREFPTR:.+]] = emitc.call "EMITC_CAST"(%[[RESPTR]]) {args = [0 : index, !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>]}
+  // CHECK-SAME:     : (!emitc.ptr<ui8>) -> !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>
+  // CHECK-NEXT: emitc.call "iree_vm_ref_move"(%[[RESREFPTR]], %arg3)
+
+  // Return ok status.
+  // CHECK-NEXT: %[[OK:.+]] = emitc.call "iree_ok_status"()
+  // CHECK-NEXT: return %[[OK]]
+  vm.import @ref_fn(%arg0 : !vm.ref<?>) -> !vm.ref<?>
+
+  vm.func @import_ref(%arg0 : !vm.ref<?>) -> !vm.ref<?> {
+    %0 = vm.call @ref_fn(%arg0) : (!vm.ref<?>) -> !vm.ref<?>
+    vm.return %0 : !vm.ref<?>
+  }
+}
+
+// -----
+
 vm.module @my_module {
 
   // Typedef structs for arguments and results
diff --git a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/func_op.mlir b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/func_op.mlir
new file mode 100644
index 0000000..e6f1418
--- /dev/null
+++ b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/func_op.mlir
@@ -0,0 +1,19 @@
+// RUN: iree-opt -split-input-file -pass-pipeline="vm.module(iree-vm-ordinal-allocation),vm.module(iree-convert-vm-to-emitc)" %s | FileCheck %s
+
+// This tests the vm.func conversion. Using the calling convention of EmitC,
+// some arguments are getting added. For more details see comments on the
+// `ConvertVMToEmitCPass` class in ConvertVMToEmitC.cpp.
+vm.module @my_module {
+  // CHECK: func @my_module_fn(%arg0: !emitc.ptr<!emitc.opaque<"iree_vm_stack_t">>
+  // CHECK-SAME:               %arg1: !emitc.ptr<!emitc.opaque<"my_module_t">>,
+  // CHECK-SAME:               %arg2: !emitc.ptr<!emitc.opaque<"my_module_state_t">>,
+  // CHECK-SAME:               %arg3: !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>,
+  // CHECK-SAME:               %arg4: i32,
+  // CHECK-SAME:               %arg5: !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>,
+  // CHECK-SAME:               %arg6: !emitc.ptr<!emitc.opaque<"int32_t">>)
+  // CHECK-SAME:               -> !emitc.opaque<"iree_status_t">
+  // CHECK-SAME: attributes {emitc.static, vm.calling_convention = "0ri_ri"}
+  vm.func @fn(%arg0 : !vm.ref<?>, %arg1 : i32) -> (!vm.ref<?>, i32) {
+    vm.return %arg0, %arg1 : !vm.ref<?>, i32
+  }
+}
diff --git a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/type_conversion.mlir b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/type_conversion.mlir
index cf9dd8f..7491f5a 100644
--- a/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/type_conversion.mlir
+++ b/iree/compiler/Dialect/VM/Conversion/VMToEmitC/test/type_conversion.mlir
@@ -39,3 +39,15 @@
     vm.return
   }
 }
+
+// -----
+
+// Test the func.func conversion, which is needed as a second step after the
+// vm.func conversion. All references in the signature should be converted to
+// emitc pointers.
+vm.module @my_module {
+  // CHECK: func @fn(%arg0: !emitc.ptr<!emitc.opaque<"iree_vm_ref_t">>, %arg1: i32)
+  func.func @fn(%arg0 : !vm.ref<?>, %arg1 : i32) -> () {
+    return
+  }
+}
diff --git a/iree/test/e2e/models/edge_detection.mlir b/iree/test/e2e/models/edge_detection.mlir
index 5a3bff0..39f66ab 100644
--- a/iree/test/e2e/models/edge_detection.mlir
+++ b/iree/test/e2e/models/edge_detection.mlir
@@ -2,7 +2,8 @@
 // RUN: iree-run-mlir --iree-input-type=mhlo -iree-hal-target-backends=dylib-llvm-aot %s -function-input="1x128x128x1xf32" | FileCheck %s
 // RUN: [[ $IREE_VULKAN_DISABLE == 1 ]] || (iree-run-mlir --iree-input-type=mhlo -iree-hal-target-backends=vulkan-spirv %s -function-input="1x128x128x1xf32" | FileCheck %s)
 
-// Image edge detection module generated by iree/colab/edge_detection.ipynb.
+// Image edge detection module generated by.
+// https://github.com/google/iree/blob/main/samples/colab/edge_detection.ipynb.
 //
 // Input : a single 128x128 pixel image as a tensor<1x128x128x1xf32>, with pixels in [0.0, 1.0]
 // Output: a single image in the same format after running edge detection
diff --git a/colab/README.md b/samples/colab/README.md
similarity index 88%
rename from colab/README.md
rename to samples/colab/README.md
index 34bbc76..5c7a53d 100644
--- a/colab/README.md
+++ b/samples/colab/README.md
@@ -7,19 +7,19 @@
 Constructs a TF module for performing image edge detection and runs it using
 IREE
 
-[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/colab/edge_detection.ipynb)
+[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/colab/edge_detection.ipynb)
 
 ### [low_level_invoke_function\.ipynb](low_level_invoke_function.ipynb)
 
 Shows off some concepts of the low level IREE python bindings
 
-[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/colab/low_level_invoke_function.ipynb)
+[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/colab/low_level_invoke_function.ipynb)
 
 ### [mnist_training\.ipynb](mnist_training.ipynb)
 
 Compile, train and execute a TensorFlow Keras neural network with IREE
 
-[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/colab/mnist_training.ipynb)
+[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/colab/mnist_training.ipynb)
 
 ### [resnet\.ipynb](resnet.ipynb)
 
@@ -27,7 +27,7 @@
 [ResNet50](https://www.tensorflow.org/api_docs/python/tf/keras/applications/ResNet50)
 model and runs it using IREE
 
-[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/colab/resnet.ipynb)
+[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/colab/resnet.ipynb)
 
 ### [tensorflow_hub_import\.ipynb](tensorflow_hub_import.ipynb)
 
@@ -35,7 +35,7 @@
 [MobileNet V2](https://tfhub.dev/google/tf2-preview/mobilenet_v2/classification)
 model, pre-processes it for import, then compiles it using IREE
 
-[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/colab/tensorflow_hub_import.ipynb)
+[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/colab/tensorflow_hub_import.ipynb)
 
 ### [tflite_text_classification\.ipynb](tflite_text_classification.ipynb)
 
@@ -43,7 +43,7 @@
 [TFLite text classification](https://www.tensorflow.org/lite/examples/text_classification/overview)
 model, and runs it using TFLite and IREE
 
-[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/colab/tflite_text_classification.ipynb)
+[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/colab/tflite_text_classification.ipynb)
 
 ## Working with GitHub
 
diff --git a/colab/dummy.py b/samples/colab/dummy.py
similarity index 100%
rename from colab/dummy.py
rename to samples/colab/dummy.py
diff --git a/colab/edge_detection.ipynb b/samples/colab/edge_detection.ipynb
similarity index 100%
rename from colab/edge_detection.ipynb
rename to samples/colab/edge_detection.ipynb
diff --git a/colab/low_level_invoke_function.ipynb b/samples/colab/low_level_invoke_function.ipynb
similarity index 100%
rename from colab/low_level_invoke_function.ipynb
rename to samples/colab/low_level_invoke_function.ipynb
diff --git a/colab/mnist_training.ipynb b/samples/colab/mnist_training.ipynb
similarity index 100%
rename from colab/mnist_training.ipynb
rename to samples/colab/mnist_training.ipynb
diff --git a/colab/resnet.ipynb b/samples/colab/resnet.ipynb
similarity index 100%
rename from colab/resnet.ipynb
rename to samples/colab/resnet.ipynb
diff --git a/colab/tensorflow_hub_import.ipynb b/samples/colab/tensorflow_hub_import.ipynb
similarity index 100%
rename from colab/tensorflow_hub_import.ipynb
rename to samples/colab/tensorflow_hub_import.ipynb
diff --git a/colab/test_notebooks.py b/samples/colab/test_notebooks.py
similarity index 96%
rename from colab/test_notebooks.py
rename to samples/colab/test_notebooks.py
index af545b8..c1ef330 100644
--- a/colab/test_notebooks.py
+++ b/samples/colab/test_notebooks.py
@@ -33,7 +33,7 @@
                                "build_tools/testing/run_python_notebook.sh")
 
     # Create a test case for each notebook in this folder.
-    notebooks_path = os.path.join(repo_root, "colab/")
+    notebooks_path = os.path.join(repo_root, "samples/colab/")
     for notebook_path in glob.glob(notebooks_path + "*.ipynb"):
       notebook_name = os.path.basename(notebook_path)
 
diff --git a/colab/tflite_text_classification.ipynb b/samples/colab/tflite_text_classification.ipynb
similarity index 100%
rename from colab/tflite_text_classification.ipynb
rename to samples/colab/tflite_text_classification.ipynb
diff --git a/samples/dynamic_shapes/README.md b/samples/dynamic_shapes/README.md
index 39f6b4b..138181d 100644
--- a/samples/dynamic_shapes/README.md
+++ b/samples/dynamic_shapes/README.md
@@ -13,7 +13,7 @@
 [`dynamic_shapes.ipynb`](./dynamic_shapes.ipynb)
 [Colab](https://research.google.com/colaboratory/) notebook:
 
-[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/iree/samples/dynamic_shapes/dynamic_shapes.ipynb)
+[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/dynamic_shapes/dynamic_shapes.ipynb)
 
 Step 3 should be performed on your development host machine
 
@@ -97,6 +97,6 @@
 5. Run the sample binary:
 
    ```
-   ../iree-build/iree/samples/dynamic_shapes/dynamic-shapes \
+   ../iree-build/samples/dynamic_shapes/dynamic-shapes \
        /path/to/dynamic_shapes_dylib.vmfb dylib
    ```
diff --git a/samples/dynamic_shapes/test.sh b/samples/dynamic_shapes/test.sh
index dacac27..7d8297f 100755
--- a/samples/dynamic_shapes/test.sh
+++ b/samples/dynamic_shapes/test.sh
@@ -18,7 +18,7 @@
 
 # 1. Run the notebook to generate `counter.mlir` and `counter_vmvx.vmfb`
 ${ROOT_DIR}/build_tools/testing/run_python_notebook.sh \
-  ${ROOT_DIR}/iree/samples/dynamic_shapes/dynamic_shapes.ipynb
+  ${ROOT_DIR}/samples/dynamic_shapes/dynamic_shapes.ipynb
 test -f ${ARTIFACTS_DIR}/dynamic_shapes.mlir && echo "dynamic_shapes.mlir exists"
 
 # 2. Build the `iree-compile` tool.
@@ -36,5 +36,5 @@
 cmake --build ${BUILD_DIR} --target iree_samples_dynamic_shapes -- -k 0
 
 # 5. Run the sample binary.
-${BUILD_DIR}/iree/samples/dynamic_shapes/dynamic-shapes \
+${BUILD_DIR}/samples/dynamic_shapes/dynamic-shapes \
   ${ARTIFACTS_DIR}/dynamic_shapes_dylib.vmfb dylib
diff --git a/samples/models/mnist.mlir b/samples/models/mnist.mlir
index 7e69a1c..b367abc 100644
--- a/samples/models/mnist.mlir
+++ b/samples/models/mnist.mlir
@@ -1,5 +1,5 @@
 // Trained MNIST model generated by
-// https://github.com/google/iree/blob/main/colab/mnist_tensorflow.ipynb.
+// https://github.com/google/iree/blob/main/samples/colab/mnist_training.ipynb.
 //
 // Model structure is from tf.keras:
 //
diff --git a/samples/variables_and_state/README.md b/samples/variables_and_state/README.md
index b03b0ee..3eda86d 100644
--- a/samples/variables_and_state/README.md
+++ b/samples/variables_and_state/README.md
@@ -13,7 +13,7 @@
 [`variables_and_state.ipynb`](./variables_and_state.ipynb)
 [Colab](https://research.google.com/colaboratory/) notebook:
 
-[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/iree/samples/variables_and_state/variables_and_state.ipynb)
+[![Open In Colab](https://colab.research.google.com/assets/colab-badge.svg)](https://colab.research.google.com/github/google/iree/blob/main/samples/variables_and_state/variables_and_state.ipynb)
 
 Steps 4-5 are in [`main.c`](./main.c)
 
@@ -73,7 +73,7 @@
 3. Run the sample binary:
 
    ```
-   ../iree-build/iree/samples/variables_and_state/variables-and-state \
+   ../iree-build/samples/variables_and_state/variables-and-state \
        /path/to/counter_vmvx.vmfb vmvx
    ```
 
@@ -101,6 +101,6 @@
 then run the program with that new VM bytecode module:
 
 ```
-../iree-build/iree/samples/variables_and_state/variables-and-state \
+../iree-build/samples/variables_and_state/variables-and-state \
     /path/to/counter_dylib.vmfb dylib
 ```
diff --git a/samples/variables_and_state/test.sh b/samples/variables_and_state/test.sh
index f78db1b..84bf523 100755
--- a/samples/variables_and_state/test.sh
+++ b/samples/variables_and_state/test.sh
@@ -18,7 +18,7 @@
 
 # 1. Run the notebook to generate `counter.mlir` and `counter_vmvx.vmfb`
 ${ROOT_DIR}/build_tools/testing/run_python_notebook.sh \
-  ${ROOT_DIR}/iree/samples/variables_and_state/variables_and_state.ipynb
+  ${ROOT_DIR}/samples/variables_and_state/variables_and_state.ipynb
 test -f ${ARTIFACTS_DIR}/counter.mlir && echo "counter.mlir exists"
 test -f ${ARTIFACTS_DIR}/counter_vmvx.vmfb && echo "counter_vmvx.vmfb exists"
 
@@ -27,5 +27,5 @@
 cmake --build ${BUILD_DIR} --target iree_samples_variables_and_state -- -k 0
 
 # 3. Run the sample binary.
-${BUILD_DIR}/iree/samples/variables_and_state/variables-and-state \
+${BUILD_DIR}/samples/variables_and_state/variables-and-state \
   ${ARTIFACTS_DIR}/counter_vmvx.vmfb vmvx
diff --git a/samples/vulkan/vulkan_inference_gui.cc b/samples/vulkan/vulkan_inference_gui.cc
index b250e55..4b1e639 100644
--- a/samples/vulkan/vulkan_inference_gui.cc
+++ b/samples/vulkan/vulkan_inference_gui.cc
@@ -26,7 +26,7 @@
 #include "iree/base/logging.h"
 
 // Compiled module embedded here to avoid file IO:
-#include "iree/samples/vulkan/simple_mul_bytecode_module_c.h"
+#include "samples/vulkan/simple_mul_bytecode_module_c.h"
 
 static VkAllocationCallbacks* g_Allocator = NULL;
 static VkInstance g_Instance = VK_NULL_HANDLE;