diff --git a/CMakeLists.txt b/CMakeLists.txt
index 36a743f..6049b1b 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -81,6 +81,7 @@
 option_if_not_defined(TINT_BUILD_SPIRV_TOOLS_FUZZER "Build SPIRV-Tools fuzzer" OFF)
 option_if_not_defined(TINT_BUILD_AST_FUZZER "Build AST fuzzer" OFF)
 option_if_not_defined(TINT_BUILD_REGEX_FUZZER "Build regex fuzzer" OFF)
+option_if_not_defined(TINT_BUILD_BENCHMARKS "Build benchmarks" OFF)
 option_if_not_defined(TINT_BUILD_TESTS "Build tests" ${TINT_BUILD_TESTS_DEFAULT})
 option_if_not_defined(TINT_BUILD_AS_OTHER_OS "Override OS detection to force building of *_other.cc files" OFF)
 option_if_not_defined(TINT_BUILD_REMOTE_COMPILE "Build the remote-compile tool for validating shaders on a remote machine" OFF)
@@ -111,6 +112,7 @@
 message(STATUS "Tint build SPIRV-Tools fuzzer: ${TINT_BUILD_SPIRV_TOOLS_FUZZER}")
 message(STATUS "Tint build AST fuzzer: ${TINT_BUILD_AST_FUZZER}")
 message(STATUS "Tint build regex fuzzer: ${TINT_BUILD_REGEX_FUZZER}")
+message(STATUS "Tint build benchmarks: ${TINT_BUILD_BENCHMARKS}")
 message(STATUS "Tint build tests: ${TINT_BUILD_TESTS}")
 message(STATUS "Tint build with ASAN: ${TINT_ENABLE_ASAN}")
 message(STATUS "Tint build with MSAN: ${TINT_ENABLE_MSAN}")
@@ -254,7 +256,7 @@
   endif(DOXYGEN_FOUND)
 endif()
 
-function(tint_default_compile_options TARGET)
+function(tint_core_compile_options TARGET)
   target_include_directories(${TARGET} PUBLIC "${TINT_ROOT_SOURCE_DIR}")
   target_include_directories(${TARGET} PUBLIC "${TINT_ROOT_SOURCE_DIR}/include")
 
@@ -263,20 +265,48 @@
         "${TINT_ROOT_SOURCE_DIR}/third_party/spirv-headers/include")
   endif()
 
-  target_compile_definitions(${TARGET} PUBLIC
-      -DTINT_BUILD_SPV_READER=$<BOOL:${TINT_BUILD_SPV_READER}>)
-  target_compile_definitions(${TARGET} PUBLIC
-      -DTINT_BUILD_WGSL_READER=$<BOOL:${TINT_BUILD_WGSL_READER}>)
-  target_compile_definitions(${TARGET} PUBLIC
-      -DTINT_BUILD_GLSL_WRITER=$<BOOL:${TINT_BUILD_GLSL_WRITER}>)
-  target_compile_definitions(${TARGET} PUBLIC
-      -DTINT_BUILD_HLSL_WRITER=$<BOOL:${TINT_BUILD_HLSL_WRITER}>)
-  target_compile_definitions(${TARGET} PUBLIC
-      -DTINT_BUILD_MSL_WRITER=$<BOOL:${TINT_BUILD_MSL_WRITER}>)
-  target_compile_definitions(${TARGET} PUBLIC
-      -DTINT_BUILD_SPV_WRITER=$<BOOL:${TINT_BUILD_SPV_WRITER}>)
-  target_compile_definitions(${TARGET} PUBLIC
-      -DTINT_BUILD_WGSL_WRITER=$<BOOL:${TINT_BUILD_WGSL_WRITER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_SPV_READER=$<BOOL:${TINT_BUILD_SPV_READER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_WGSL_READER=$<BOOL:${TINT_BUILD_WGSL_READER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_GLSL_WRITER=$<BOOL:${TINT_BUILD_GLSL_WRITER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_HLSL_WRITER=$<BOOL:${TINT_BUILD_HLSL_WRITER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_MSL_WRITER=$<BOOL:${TINT_BUILD_MSL_WRITER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_SPV_WRITER=$<BOOL:${TINT_BUILD_SPV_WRITER}>)
+  target_compile_definitions(${TARGET} PUBLIC -DTINT_BUILD_WGSL_WRITER=$<BOOL:${TINT_BUILD_WGSL_WRITER}>)
+
+  if (COMPILER_IS_LIKE_GNU)
+    target_compile_options(${TARGET} PRIVATE
+      -std=c++17
+      -fno-exceptions
+      -fno-rtti
+    )
+
+    if (${TINT_ENABLE_MSAN})
+      target_compile_options(${TARGET} PRIVATE -fsanitize=memory)
+      target_link_options(${TARGET} PRIVATE -fsanitize=memory)
+    elseif (${TINT_ENABLE_ASAN})
+      target_compile_options(${TARGET} PRIVATE -fsanitize=address)
+      target_link_options(${TARGET} PRIVATE -fsanitize=address)
+    elseif (${TINT_ENABLE_UBSAN})
+      target_compile_options(${TARGET} PRIVATE -fsanitize=undefined)
+      target_link_options(${TARGET} PRIVATE -fsanitize=undefined)
+    endif()
+  endif(COMPILER_IS_LIKE_GNU)
+
+  if (TINT_EMIT_COVERAGE)
+    if(CMAKE_CXX_COMPILER_ID MATCHES "GNU")
+        target_compile_options(${TARGET} PRIVATE "--coverage")
+        target_link_options(${TARGET} PRIVATE "gcov")
+    elseif(CMAKE_CXX_COMPILER_ID MATCHES "Clang")
+        target_compile_options(${TARGET} PRIVATE "-fprofile-instr-generate" "-fcoverage-mapping")
+        target_link_options(${TARGET} PRIVATE "-fprofile-instr-generate" "-fcoverage-mapping")
+    else()
+        message(FATAL_ERROR "Coverage generation not supported for the ${CMAKE_CXX_COMPILER_ID} toolchain")
+    endif()
+  endif(TINT_EMIT_COVERAGE)
+endfunction()
+
+function(tint_default_compile_options TARGET)
+  tint_core_compile_options(${TARGET})
 
   set(COMMON_GNU_OPTIONS
     -Wall
@@ -299,11 +329,8 @@
     -Weverything
   )
 
-  if (${COMPILER_IS_LIKE_GNU})
+  if (COMPILER_IS_LIKE_GNU)
     target_compile_options(${TARGET} PRIVATE
-      -std=c++17
-      -fno-exceptions
-      -fno-rtti
       -pedantic-errors
       ${COMMON_GNU_OPTIONS}
     )
@@ -314,30 +341,7 @@
         ${COMMON_CLANG_OPTIONS}
       )
     endif()
-
-    if (${TINT_ENABLE_MSAN})
-      target_compile_options(${TARGET} PRIVATE -fsanitize=memory)
-      target_link_options(${TARGET} PRIVATE -fsanitize=memory)
-    elseif (${TINT_ENABLE_ASAN})
-      target_compile_options(${TARGET} PRIVATE -fsanitize=address)
-      target_link_options(${TARGET} PRIVATE -fsanitize=address)
-    elseif (${TINT_ENABLE_UBSAN})
-      target_compile_options(${TARGET} PRIVATE -fsanitize=undefined)
-      target_link_options(${TARGET} PRIVATE -fsanitize=undefined)
-    endif()
-  endif()
-
-  if (${TINT_EMIT_COVERAGE})
-    if(CMAKE_CXX_COMPILER_ID MATCHES "GNU")
-        target_compile_options(${TARGET} PRIVATE "--coverage")
-        target_link_options(${TARGET} PRIVATE "gcov")
-    elseif(CMAKE_CXX_COMPILER_ID MATCHES "Clang")
-        target_compile_options(${TARGET} PRIVATE "-fprofile-instr-generate" "-fcoverage-mapping")
-        target_link_options(${TARGET} PRIVATE "-fprofile-instr-generate" "-fcoverage-mapping")
-    else()
-        message(FATAL_ERROR "Coverage generation not supported for the ${CMAKE_CXX_COMPILER_ID} toolchain")
-    endif()
-  endif()
+  endif(COMPILER_IS_LIKE_GNU)
 
   if (MSVC)
     # Specify /EHs for exception handling.
@@ -379,7 +383,6 @@
       )
     endif()
   endif()
-
 endfunction()
 
 add_subdirectory(third_party)
diff --git a/DEPS b/DEPS
index 6bb6635..fbd2af5 100644
--- a/DEPS
+++ b/DEPS
@@ -10,6 +10,7 @@
   'chromium_git':  'https://chromium.googlesource.com',
   'github': '/external/github.com',
 
+  'benchmark_revision': 'e991355c02b93fe17713efe04cbc2e278e00fdbd',
   'build_revision': '555c8b467c21e2c4b22d00e87e3faa0431df9ac2',
   'buildtools_revision': 'f78b4b9f33bd8ef9944d5ce643daff1c31880189',
   'catapult_revision': 'fa35beefb3429605035f98211ddb8750dee6a13d',
@@ -101,6 +102,9 @@
   'third_party/catapult': Var('chromium_git') + '/catapult.git@' +
       Var('catapult_revision'),
 
+  'third_party/benchmark': Var('chromium_git') + Var('github') +
+      '/google/benchmark.git@' + Var('benchmark_revision'),
+
   'third_party/googletest': Var('chromium_git') + Var('github') +
       '/google/googletest.git@' + Var('googletest_revision'),
 
diff --git a/fuzzers/tint_black_box_fuzz_target.cc b/fuzzers/tint_black_box_fuzz_target.cc
index cefaa76..f2f2f68 100644
--- a/fuzzers/tint_black_box_fuzz_target.cc
+++ b/fuzzers/tint_black_box_fuzz_target.cc
@@ -56,13 +56,7 @@
   }
 
   fseek(file, 0, SEEK_END);
-  auto tell_file_size = ftell(file);
-  if (tell_file_size <= 0) {
-    std::cerr << "Input file of incorrect size: " << input_file << std::endl;
-    fclose(file);
-    return {};
-  }
-  const auto file_size = static_cast<size_t>(tell_file_size);
+  const auto file_size = static_cast<size_t>(ftell(file));
   if (0 != (file_size % sizeof(T))) {
     std::cerr << "File " << input_file
               << " does not contain an integral number of objects: "
diff --git a/kokoro/linux/docker.sh b/kokoro/linux/docker.sh
index cef68d1..9818060 100755
--- a/kokoro/linux/docker.sh
+++ b/kokoro/linux/docker.sh
@@ -84,7 +84,8 @@
 
     COMMON_CMAKE_FLAGS=""
     COMMON_CMAKE_FLAGS+=" -DCMAKE_BUILD_TYPE=${BUILD_TYPE}"
-    COMMON_CMAKE_FLAGS+=" -DTINT_DOCS_WARN_AS_ERROR=ON"
+    COMMON_CMAKE_FLAGS+=" -DTINT_DOCS_WARN_AS_ERROR=1"
+    COMMON_CMAKE_FLAGS+=" -DTINT_BUILD_BENCHMARKS=1"
 
     if [ "$BUILD_TOOLCHAIN" == "clang" ]; then
         using clang-10.0.0
@@ -155,9 +156,9 @@
 
     status "Checking disabling all readers and writers also builds"
     show_cmds
-        cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS} -DTINT_BUILD_SPV_READER=OFF -DTINT_BUILD_SPV_WRITER=OFF -DTINT_BUILD_WGSL_READER=OFF -DTINT_BUILD_WGSL_WRITER=OFF -DTINT_BUILD_MSL_WRITER=OFF -DTINT_BUILD_HLSL_WRITER=OFF
+        cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS} -DTINT_BUILD_SPV_READER=OFF -DTINT_BUILD_SPV_WRITER=OFF -DTINT_BUILD_WGSL_READER=OFF -DTINT_BUILD_WGSL_WRITER=OFF -DTINT_BUILD_MSL_WRITER=OFF -DTINT_BUILD_HLSL_WRITER=OFF -DTINT_BUILD_BENCHMARKS=OFF
         cmake --build . -- --jobs=$(nproc)
-        cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS} -DTINT_BUILD_SPV_READER=ON -DTINT_BUILD_SPV_WRITER=ON -DTINT_BUILD_WGSL_READER=ON -DTINT_BUILD_WGSL_WRITER=ON -DTINT_BUILD_MSL_WRITER=ON -DTINT_BUILD_HLSL_WRITER=ON
+        cmake ${SRC_DIR} ${CMAKE_FLAGS} ${COMMON_CMAKE_FLAGS} -DTINT_BUILD_SPV_READER=ON -DTINT_BUILD_SPV_WRITER=ON -DTINT_BUILD_WGSL_READER=ON -DTINT_BUILD_WGSL_WRITER=ON -DTINT_BUILD_MSL_WRITER=ON -DTINT_BUILD_HLSL_WRITER=ON -DTINT_BUILD_BENCHMARKS=ON
     hide_cmds
 else
     status "Unsupported build system: $BUILD_SYSTEM"
diff --git a/kokoro/windows/build.bat b/kokoro/windows/build.bat
index a7d28ce..e22d5d9 100644
--- a/kokoro/windows/build.bat
+++ b/kokoro/windows/build.bat
@@ -116,7 +116,7 @@
 @echo on
 mkdir %BUILD_DIR%
 cd /d %BUILD_DIR%
-set COMMON_CMAKE_FLAGS=-DTINT_BUILD_DOCS=O -DCMAKE_BUILD_TYPE=%BUILD_TYPE%
+set COMMON_CMAKE_FLAGS=-DTINT_BUILD_DOCS=O -DTINT_BUILD_BENCHMARKS=1 -DCMAKE_BUILD_TYPE=%BUILD_TYPE%
 @echo off
 
 call :status "Building tint"
diff --git a/samples/main.cc b/samples/main.cc
index 4348636..87f0b54 100644
--- a/samples/main.cc
+++ b/samples/main.cc
@@ -437,13 +437,7 @@
   }
 
   fseek(file, 0, SEEK_END);
-  uint64_t tell_file_size = static_cast<uint64_t>(ftell(file));
-  if (tell_file_size <= 0) {
-    std::cerr << "Input file of incorrect size: " << input_file << std::endl;
-    fclose(file);
-    return {};
-  }
-  const auto file_size = static_cast<size_t>(tell_file_size);
+  const auto file_size = static_cast<size_t>(ftell(file));
   if (0 != (file_size % sizeof(T))) {
     std::cerr << "File " << input_file
               << " does not contain an integral number of objects: "
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index be42ae5..d760dcb 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -598,7 +598,10 @@
   endif()
 endif()
 
-if(${TINT_BUILD_TESTS})
+################################################################################
+# Tests
+################################################################################
+if(TINT_BUILD_TESTS)
   set(TINT_TEST_SRCS
     ast/alias_test.cc
     ast/array_test.cc
@@ -1114,6 +1117,7 @@
   endif()
 
   add_executable(tint_unittests ${TINT_TEST_SRCS})
+  set_target_properties(${target} PROPERTIES FOLDER "Tests")
 
   if(NOT MSVC)
     target_compile_options(tint_unittests PRIVATE
@@ -1133,4 +1137,41 @@
   endif()
 
   add_test(NAME tint_unittests COMMAND tint_unittests)
-endif()
+endif(TINT_BUILD_TESTS)
+
+################################################################################
+# Benchmarks
+################################################################################
+if(TINT_BUILD_BENCHMARKS)
+  if(NOT TINT_BUILD_WGSL_READER)
+    message(FATAL_ERROR "TINT_BUILD_BENCHMARKS requires TINT_BUILD_WGSL_READER")
+  endif()
+
+  set(TINT_BENCHMARK_SRC
+    "benchmark/benchmark.cc"
+    "reader/wgsl/parser_bench.cc"
+  )
+
+  if (${TINT_BUILD_GLSL_WRITER})
+    list(APPEND TINT_BENCHMARK_SRC writer/glsl/generator_bench.cc)
+  endif()
+  if (${TINT_BUILD_HLSL_WRITER})
+    list(APPEND TINT_BENCHMARK_SRC writer/hlsl/generator_bench.cc)
+  endif()
+  if (${TINT_BUILD_MSL_WRITER})
+    list(APPEND TINT_BENCHMARK_SRC writer/msl/generator_bench.cc)
+  endif()
+  if (${TINT_BUILD_SPV_WRITER})
+    list(APPEND TINT_BENCHMARK_SRC writer/spirv/generator_bench.cc)
+  endif()
+  if (${TINT_BUILD_WGSL_WRITER})
+    list(APPEND TINT_BENCHMARK_SRC writer/wgsl/generator_bench.cc)
+  endif()
+
+  add_executable(tint-benchmark ${TINT_BENCHMARK_SRC})
+  set_target_properties(${target} PROPERTIES FOLDER "Benchmarks")
+
+  tint_core_compile_options(tint-benchmark)
+
+  target_link_libraries(tint-benchmark PRIVATE benchmark::benchmark libtint)
+endif(TINT_BUILD_BENCHMARKS)
diff --git a/src/benchmark/benchmark.cc b/src/benchmark/benchmark.cc
new file mode 100644
index 0000000..813af84
--- /dev/null
+++ b/src/benchmark/benchmark.cc
@@ -0,0 +1,123 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/benchmark/benchmark.h"
+
+#include <filesystem>
+#include <sstream>
+#include <utility>
+#include <vector>
+
+namespace tint::benchmark {
+namespace {
+
+std::filesystem::path kInputFileDir;
+
+/// Copies the content from the file named `input_file` to `buffer`,
+/// assuming each element in the file is of type `T`.  If any error occurs,
+/// writes error messages to the standard error stream and returns false.
+/// Assumes the size of a `T` object is divisible by its required alignment.
+/// @returns true if we successfully read the file.
+template <typename T>
+std::variant<std::vector<T>, Error> ReadFile(const std::string& input_file) {
+  FILE* file = nullptr;
+#if defined(_MSC_VER)
+  fopen_s(&file, input_file.c_str(), "rb");
+#else
+  file = fopen(input_file.c_str(), "rb");
+#endif
+  if (!file) {
+    return Error{"Failed to open " + input_file};
+  }
+
+  fseek(file, 0, SEEK_END);
+  const auto file_size = static_cast<size_t>(ftell(file));
+  if (0 != (file_size % sizeof(T))) {
+    std::stringstream err;
+    err << "File " << input_file
+        << " does not contain an integral number of objects: " << file_size
+        << " bytes in the file, require " << sizeof(T) << " bytes per object";
+    fclose(file);
+    return Error{err.str()};
+  }
+  fseek(file, 0, SEEK_SET);
+
+  std::vector<T> buffer;
+  buffer.resize(file_size / sizeof(T));
+
+  size_t bytes_read = fread(buffer.data(), 1, file_size, file);
+  fclose(file);
+  if (bytes_read != file_size) {
+    return Error{"Failed to read " + input_file};
+  }
+
+  return buffer;
+}
+
+bool FindBenchmarkInputDir() {
+  // Attempt to find the benchmark input files by searching up from the current
+  // working directory.
+  auto path = std::filesystem::current_path();
+  while (std::filesystem::is_directory(path)) {
+    auto test = path / "test" / "benchmark";
+    if (std::filesystem::is_directory(test)) {
+      kInputFileDir = test;
+      return true;
+    }
+    auto parent = path.parent_path();
+    if (path == parent) {
+      break;
+    }
+    path = parent;
+  }
+  return false;
+}
+
+}  // namespace
+
+std::variant<tint::Source::File, Error> LoadInputFile(std::string name) {
+  auto path = (kInputFileDir / name).string();
+  auto data = ReadFile<uint8_t>(path);
+  if (auto* buf = std::get_if<std::vector<uint8_t>>(&data)) {
+    return tint::Source::File(path, std::string(buf->begin(), buf->end()));
+  }
+  return std::get<Error>(data);
+}
+
+std::variant<ProgramAndFile, Error> LoadProgram(std::string name) {
+  auto res = benchmark::LoadInputFile(name);
+  if (auto err = std::get_if<benchmark::Error>(&res)) {
+    return *err;
+  }
+  auto& file = std::get<Source::File>(res);
+  auto program = reader::wgsl::Parse(&file);
+  if (program.Diagnostics().contains_errors()) {
+    return Error{program.Diagnostics().str()};
+  }
+  return ProgramAndFile{std::move(program), std::move(file)};
+}
+
+}  // namespace tint::benchmark
+
+int main(int argc, char** argv) {
+  ::benchmark::Initialize(&argc, argv);
+  if (::benchmark::ReportUnrecognizedArguments(argc, argv)) {
+    return 1;
+  }
+  if (!tint::benchmark::FindBenchmarkInputDir()) {
+    std::cerr << "failed to locate benchmark input files" << std::endl;
+    return 1;
+  }
+  ::benchmark::RunSpecifiedBenchmarks();
+}
diff --git a/src/benchmark/benchmark.h b/src/benchmark/benchmark.h
new file mode 100644
index 0000000..40de31d
--- /dev/null
+++ b/src/benchmark/benchmark.h
@@ -0,0 +1,69 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_BENCHMARK_BENCHMARK_H_
+#define SRC_BENCHMARK_BENCHMARK_H_
+
+#include <memory>
+#include <string>
+#include <variant>  // NOLINT: Found C system header after C++ system header.
+
+#include "benchmark/benchmark.h"
+#include "src/utils/concat.h"
+#include "tint/tint.h"
+
+namespace tint::benchmark {
+
+/// Error indicates an operation did not complete successfully.
+struct Error {
+  /// The error message.
+  std::string msg;
+};
+
+/// ProgramAndFile holds a Program and a Source::File.
+struct ProgramAndFile {
+  /// The tint program parsed from file.
+  Program program;
+  /// The source file
+  Source::File file;
+};
+
+/// LoadInputFile attempts to load a benchmark input file with the given file
+/// name.
+/// @param name the file name
+/// @returns either the loaded Source::File or an Error
+std::variant<Source::File, Error> LoadInputFile(std::string name);
+
+/// LoadInputFile attempts to load a benchmark input program with the given file
+/// name.
+/// @param name the file name
+/// @returns either the loaded Program or an Error
+std::variant<ProgramAndFile, Error> LoadProgram(std::string name);
+
+/// Declares a benchmark with the given function and WGSL file name
+#define TINT_BENCHMARK_WGSL_PROGRAM(FUNC, WGSL_NAME) \
+  BENCHMARK_CAPTURE(FUNC, WGSL_NAME, WGSL_NAME);
+
+/// Declares a set of benchmarks for the given function using a list of WGSL
+/// files in `<tint>/test/benchmark`.
+#define TINT_BENCHMARK_WGSL_PROGRAMS(FUNC)                   \
+  TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "empty.wgsl");           \
+  TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "particles.wgsl");       \
+  TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "simple_fragment.wgsl"); \
+  TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "simple_vertex.wgsl");   \
+  TINT_BENCHMARK_WGSL_PROGRAM(FUNC, "simple_compute.wgsl");
+
+}  // namespace tint::benchmark
+
+#endif  // SRC_BENCHMARK_BENCHMARK_H_
diff --git a/src/reader/wgsl/parser_bench.cc b/src/reader/wgsl/parser_bench.cc
new file mode 100644
index 0000000..95fa446
--- /dev/null
+++ b/src/reader/wgsl/parser_bench.cc
@@ -0,0 +1,40 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "src/benchmark/benchmark.h"
+
+namespace tint::reader::wgsl {
+namespace {
+
+void ParseWGSL(::benchmark::State& state, std::string input_name) {
+  auto res = benchmark::LoadInputFile(input_name);
+  if (auto err = std::get_if<benchmark::Error>(&res)) {
+    state.SkipWithError(err->msg.c_str());
+    return;
+  }
+  auto& file = std::get<Source::File>(res);
+  for (auto _ : state) {
+    auto res = Parse(&file);
+    if (res.Diagnostics().contains_errors()) {
+      state.SkipWithError(res.Diagnostics().str().c_str());
+    }
+  }
+}
+
+TINT_BENCHMARK_WGSL_PROGRAMS(ParseWGSL);
+
+}  // namespace
+}  // namespace tint::reader::wgsl
diff --git a/src/writer/glsl/generator_bench.cc b/src/writer/glsl/generator_bench.cc
new file mode 100644
index 0000000..2ac109f
--- /dev/null
+++ b/src/writer/glsl/generator_bench.cc
@@ -0,0 +1,50 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "src/ast/module.h"
+#include "src/benchmark/benchmark.h"
+
+namespace tint::writer::glsl {
+namespace {
+
+void GenerateGLSL(::benchmark::State& state, std::string input_name) {
+  auto res = benchmark::LoadProgram(input_name);
+  if (auto err = std::get_if<benchmark::Error>(&res)) {
+    state.SkipWithError(err->msg.c_str());
+    return;
+  }
+  auto& program = std::get<benchmark::ProgramAndFile>(res).program;
+  std::vector<std::string> entry_points;
+  for (auto& fn : program.AST().Functions()) {
+    if (fn->IsEntryPoint()) {
+      entry_points.emplace_back(program.Symbols().NameFor(fn->symbol));
+    }
+  }
+
+  for (auto _ : state) {
+    for (auto& ep : entry_points) {
+      auto res = Generate(&program, {}, ep);
+      if (!res.error.empty()) {
+        state.SkipWithError(res.error.c_str());
+      }
+    }
+  }
+}
+
+TINT_BENCHMARK_WGSL_PROGRAMS(GenerateGLSL);
+
+}  // namespace
+}  // namespace tint::writer::glsl
diff --git a/src/writer/hlsl/generator_bench.cc b/src/writer/hlsl/generator_bench.cc
new file mode 100644
index 0000000..e9168af
--- /dev/null
+++ b/src/writer/hlsl/generator_bench.cc
@@ -0,0 +1,40 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "src/benchmark/benchmark.h"
+
+namespace tint::writer::hlsl {
+namespace {
+
+void GenerateHLSL(::benchmark::State& state, std::string input_name) {
+  auto res = benchmark::LoadProgram(input_name);
+  if (auto err = std::get_if<benchmark::Error>(&res)) {
+    state.SkipWithError(err->msg.c_str());
+    return;
+  }
+  auto& program = std::get<benchmark::ProgramAndFile>(res).program;
+  for (auto _ : state) {
+    auto res = Generate(&program, {});
+    if (!res.error.empty()) {
+      state.SkipWithError(res.error.c_str());
+    }
+  }
+}
+
+TINT_BENCHMARK_WGSL_PROGRAMS(GenerateHLSL);
+
+}  // namespace
+}  // namespace tint::writer::hlsl
diff --git a/src/writer/msl/generator_bench.cc b/src/writer/msl/generator_bench.cc
new file mode 100644
index 0000000..d2b53df
--- /dev/null
+++ b/src/writer/msl/generator_bench.cc
@@ -0,0 +1,40 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "src/benchmark/benchmark.h"
+
+namespace tint::writer::msl {
+namespace {
+
+void GenerateMSL(::benchmark::State& state, std::string input_name) {
+  auto res = benchmark::LoadProgram(input_name);
+  if (auto err = std::get_if<benchmark::Error>(&res)) {
+    state.SkipWithError(err->msg.c_str());
+    return;
+  }
+  auto& program = std::get<benchmark::ProgramAndFile>(res).program;
+  for (auto _ : state) {
+    auto res = Generate(&program, {});
+    if (!res.error.empty()) {
+      state.SkipWithError(res.error.c_str());
+    }
+  }
+}
+
+TINT_BENCHMARK_WGSL_PROGRAMS(GenerateMSL);
+
+}  // namespace
+}  // namespace tint::writer::msl
diff --git a/src/writer/spirv/generator_bench.cc b/src/writer/spirv/generator_bench.cc
new file mode 100644
index 0000000..fb937d1
--- /dev/null
+++ b/src/writer/spirv/generator_bench.cc
@@ -0,0 +1,40 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "src/benchmark/benchmark.h"
+
+namespace tint::writer::spirv {
+namespace {
+
+void GenerateSPIRV(::benchmark::State& state, std::string input_name) {
+  auto res = benchmark::LoadProgram(input_name);
+  if (auto err = std::get_if<benchmark::Error>(&res)) {
+    state.SkipWithError(err->msg.c_str());
+    return;
+  }
+  auto& program = std::get<benchmark::ProgramAndFile>(res).program;
+  for (auto _ : state) {
+    auto res = Generate(&program, {});
+    if (!res.error.empty()) {
+      state.SkipWithError(res.error.c_str());
+    }
+  }
+}
+
+TINT_BENCHMARK_WGSL_PROGRAMS(GenerateSPIRV);
+
+}  // namespace
+}  // namespace tint::writer::spirv
diff --git a/src/writer/wgsl/generator_bench.cc b/src/writer/wgsl/generator_bench.cc
new file mode 100644
index 0000000..524d9cf
--- /dev/null
+++ b/src/writer/wgsl/generator_bench.cc
@@ -0,0 +1,40 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "src/benchmark/benchmark.h"
+
+namespace tint::writer::wgsl {
+namespace {
+
+void GenerateWGSL(::benchmark::State& state, std::string input_name) {
+  auto res = benchmark::LoadProgram(input_name);
+  if (auto err = std::get_if<benchmark::Error>(&res)) {
+    state.SkipWithError(err->msg.c_str());
+    return;
+  }
+  auto& program = std::get<benchmark::ProgramAndFile>(res).program;
+  for (auto _ : state) {
+    auto res = Generate(&program, {});
+    if (!res.error.empty()) {
+      state.SkipWithError(res.error.c_str());
+    }
+  }
+}
+
+TINT_BENCHMARK_WGSL_PROGRAMS(GenerateWGSL);
+
+}  // namespace
+}  // namespace tint::writer::wgsl
diff --git a/test/benchmark/empty.wgsl b/test/benchmark/empty.wgsl
new file mode 100644
index 0000000..e69de29
--- /dev/null
+++ b/test/benchmark/empty.wgsl
diff --git a/test/benchmark/empty.wgsl.expected.hlsl b/test/benchmark/empty.wgsl.expected.hlsl
new file mode 100644
index 0000000..da1baa5
--- /dev/null
+++ b/test/benchmark/empty.wgsl.expected.hlsl
@@ -0,0 +1,4 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
diff --git a/test/benchmark/empty.wgsl.expected.msl b/test/benchmark/empty.wgsl.expected.msl
new file mode 100644
index 0000000..466ceaa
--- /dev/null
+++ b/test/benchmark/empty.wgsl.expected.msl
@@ -0,0 +1,3 @@
+#include <metal_stdlib>
+
+using namespace metal;
diff --git a/test/benchmark/empty.wgsl.expected.spvasm b/test/benchmark/empty.wgsl.expected.spvasm
new file mode 100644
index 0000000..65bef94
--- /dev/null
+++ b/test/benchmark/empty.wgsl.expected.spvasm
@@ -0,0 +1,16 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 5
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test/benchmark/empty.wgsl.expected.wgsl b/test/benchmark/empty.wgsl.expected.wgsl
new file mode 100644
index 0000000..e69de29
--- /dev/null
+++ b/test/benchmark/empty.wgsl.expected.wgsl
diff --git a/test/benchmark/particles.wgsl b/test/benchmark/particles.wgsl
new file mode 100644
index 0000000..4734a37
--- /dev/null
+++ b/test/benchmark/particles.wgsl
@@ -0,0 +1,183 @@
+////////////////////////////////////////////////////////////////////////////////
+// Utilities
+////////////////////////////////////////////////////////////////////////////////
+var<private> rand_seed : vec2<f32>;
+
+fn rand() -> f32 {
+    rand_seed.x = fract(cos(dot(rand_seed, vec2<f32>(23.14077926, 232.61690225))) * 136.8168);
+    rand_seed.y = fract(cos(dot(rand_seed, vec2<f32>(54.47856553, 345.84153136))) * 534.7645);
+    return rand_seed.y;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Vertex shader
+////////////////////////////////////////////////////////////////////////////////
+struct RenderParams {
+  modelViewProjectionMatrix : mat4x4<f32>;
+  right : vec3<f32>;
+  up    : vec3<f32>;
+};
+[[binding(0), group(0)]] var<uniform> render_params : RenderParams;
+
+struct VertexInput {
+  [[location(0)]] position : vec3<f32>;
+  [[location(1)]] color    : vec4<f32>;
+  [[location(2)]] quad_pos : vec2<f32>; // -1..+1
+};
+
+struct VertexOutput {
+  [[builtin(position)]] position : vec4<f32>;
+  [[location(0)]]       color    : vec4<f32>;
+  [[location(1)]]       quad_pos : vec2<f32>; // -1..+1
+};
+
+[[stage(vertex)]]
+fn vs_main(in : VertexInput) -> VertexOutput {
+  var quad_pos = mat2x3<f32>(render_params.right, render_params.up) * in.quad_pos;
+  var position = in.position + quad_pos * 0.01;
+  var out : VertexOutput;
+  out.position = render_params.modelViewProjectionMatrix * vec4<f32>(position, 1.0);
+  out.color = in.color;
+  out.quad_pos = in.quad_pos;
+  return out;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Fragment shader
+////////////////////////////////////////////////////////////////////////////////
+[[stage(fragment)]]
+fn fs_main(in : VertexOutput) -> [[location(0)]] vec4<f32> {
+  var color = in.color;
+  // Apply a circular particle alpha mask
+  color.a = color.a * max(1.0 - length(in.quad_pos), 0.0);
+  return color;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Simulation Compute shader
+////////////////////////////////////////////////////////////////////////////////
+struct SimulationParams {
+  deltaTime : f32;
+  seed : vec4<f32>;
+};
+
+struct Particle {
+  position : vec3<f32>;
+  lifetime : f32;
+  color    : vec4<f32>;
+  velocity : vec3<f32>;
+};
+
+struct Particles {
+  particles : array<Particle>;
+};
+
+[[binding(0), group(0)]] var<uniform> sim_params : SimulationParams;
+[[binding(1), group(0)]] var<storage, read_write> data : Particles;
+[[binding(2), group(0)]] var texture : texture_2d<f32>;
+
+[[stage(compute), workgroup_size(64)]]
+fn simulate([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+  rand_seed = (sim_params.seed.xy + vec2<f32>(GlobalInvocationID.xy)) * sim_params.seed.zw;
+
+  let idx = GlobalInvocationID.x;
+  var particle = data.particles[idx];
+
+  // Apply gravity
+  particle.velocity.z = particle.velocity.z - sim_params.deltaTime * 0.5;
+
+  // Basic velocity integration
+  particle.position = particle.position + sim_params.deltaTime * particle.velocity;
+
+  // Age each particle. Fade out before vanishing.
+  particle.lifetime = particle.lifetime - sim_params.deltaTime;
+  particle.color.a = smoothStep(0.0, 0.5, particle.lifetime);
+
+  // If the lifetime has gone negative, then the particle is dead and should be
+  // respawned.
+  if (particle.lifetime < 0.0) {
+    // Use the probability map to find where the particle should be spawned.
+    // Starting with the 1x1 mip level.
+    var coord = vec2<i32>(0, 0);
+    for (var level = textureNumLevels(texture) - 1; level > 0; level = level - 1) {
+      // Load the probability value from the mip-level
+      // Generate a random number and using the probabilty values, pick the
+      // next texel in the next largest mip level:
+      //
+      // 0.0    probabilites.r    probabilites.g    probabilites.b   1.0
+      //  |              |              |              |              |
+      //  |   TOP-LEFT   |  TOP-RIGHT   | BOTTOM-LEFT  | BOTTOM_RIGHT |
+      //
+      let probabilites = textureLoad(texture, coord, level);
+      let value = vec4<f32>(rand());
+      let mask = (value >= vec4<f32>(0.0, probabilites.xyz)) & (value < probabilites);
+      coord = coord * 2;
+      coord.x = coord.x + select(0, 1, any(mask.yw)); // x  y
+      coord.y = coord.y + select(0, 1, any(mask.zw)); // z  w
+    }
+    let uv = vec2<f32>(coord) / vec2<f32>(textureDimensions(texture));
+    particle.position = vec3<f32>((uv - 0.5) * 3.0 * vec2<f32>(1.0, -1.0), 0.0);
+    particle.color = textureLoad(texture, coord, 0);
+    particle.velocity.x = (rand() - 0.5) * 0.1;
+    particle.velocity.y = (rand() - 0.5) * 0.1;
+    particle.velocity.z = rand() * 0.3;
+    particle.lifetime = 0.5 + rand() * 2.0;
+  }
+
+  // Store the new particle value
+  data.particles[idx] = particle;
+}
+
+struct UBO {
+  width : u32;
+};
+
+struct Buffer {
+  weights : array<f32>;
+};
+
+[[binding(3), group(0)]] var<uniform> ubo : UBO;
+[[binding(4), group(0)]] var<storage, read> buf_in : Buffer;
+[[binding(5), group(0)]] var<storage, read_write> buf_out : Buffer;
+[[binding(6), group(0)]] var tex_in : texture_2d<f32>;
+[[binding(7), group(0)]] var tex_out : texture_storage_2d<rgba8unorm, write>;
+
+////////////////////////////////////////////////////////////////////////////////
+// import_level
+//
+// Loads the alpha channel from a texel of the source image, and writes it to
+// the buf_out.weights.
+////////////////////////////////////////////////////////////////////////////////
+[[stage(compute), workgroup_size(64)]]
+fn import_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+  _ = &buf_in;
+  let offset = coord.x + coord.y * ubo.width;
+  buf_out.weights[offset] = textureLoad(tex_in, vec2<i32>(coord.xy), 0).w;
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// export_level
+//
+// Loads 4 f32 weight values from buf_in.weights, and stores summed value into
+// buf_out.weights, along with the calculated 'probabilty' vec4 values into the
+// mip level of tex_out. See simulate() in particle.wgsl to understand the
+// probability logic.
+////////////////////////////////////////////////////////////////////////////////
+[[stage(compute), workgroup_size(64)]]
+fn export_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+  if (all(coord.xy < vec2<u32>(textureDimensions(tex_out)))) {
+    let dst_offset = coord.x    + coord.y    * ubo.width;
+    let src_offset = coord.x*2u + coord.y*2u * ubo.width;
+
+    let a = buf_in.weights[src_offset + 0u];
+    let b = buf_in.weights[src_offset + 1u];
+    let c = buf_in.weights[src_offset + 0u + ubo.width];
+    let d = buf_in.weights[src_offset + 1u + ubo.width];
+    let sum = dot(vec4<f32>(a, b, c, d), vec4<f32>(1.0));
+
+    buf_out.weights[dst_offset] = sum / 4.0;
+
+    let probabilities = vec4<f32>(a, a+b, a+b+c, sum) / max(sum, 0.0001);
+    textureStore(tex_out, vec2<i32>(coord.xy), probabilities);
+  }
+}
diff --git a/test/benchmark/particles.wgsl.expected.hlsl b/test/benchmark/particles.wgsl.expected.hlsl
new file mode 100644
index 0000000..2a91f43
--- /dev/null
+++ b/test/benchmark/particles.wgsl.expected.hlsl
@@ -0,0 +1,204 @@
+static float2 rand_seed = float2(0.0f, 0.0f);
+
+float rand() {
+  rand_seed.x = frac((cos(dot(rand_seed, float2(23.140779495f, 232.616897583f))) * 136.816802979f));
+  rand_seed.y = frac((cos(dot(rand_seed, float2(54.478565216f, 345.841522217f))) * 534.764526367f));
+  return rand_seed.y;
+}
+
+cbuffer cbuffer_render_params : register(b0, space0) {
+  uint4 render_params[6];
+};
+
+struct VertexInput {
+  float3 position;
+  float4 color;
+  float2 quad_pos;
+};
+struct VertexOutput {
+  float4 position;
+  float4 color;
+  float2 quad_pos;
+};
+struct tint_symbol_5 {
+  float3 position : TEXCOORD0;
+  float4 color : TEXCOORD1;
+  float2 quad_pos : TEXCOORD2;
+};
+struct tint_symbol_6 {
+  float4 color : TEXCOORD0;
+  float2 quad_pos : TEXCOORD1;
+  float4 position : SV_Position;
+};
+
+float4x4 tint_symbol_17(uint4 buffer[6], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  const uint scalar_offset_2 = ((offset + 32u)) / 4;
+  const uint scalar_offset_3 = ((offset + 48u)) / 4;
+  return float4x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4]));
+}
+
+VertexOutput vs_main_inner(VertexInput tint_symbol) {
+  float3 quad_pos = mul(tint_symbol.quad_pos, float2x3(asfloat(render_params[4].xyz), asfloat(render_params[5].xyz)));
+  float3 position = (tint_symbol.position + (quad_pos * 0.01f));
+  VertexOutput tint_symbol_1 = (VertexOutput)0;
+  tint_symbol_1.position = mul(float4(position, 1.0f), tint_symbol_17(render_params, 0u));
+  tint_symbol_1.color = tint_symbol.color;
+  tint_symbol_1.quad_pos = tint_symbol.quad_pos;
+  return tint_symbol_1;
+}
+
+tint_symbol_6 vs_main(tint_symbol_5 tint_symbol_4) {
+  const VertexInput tint_symbol_32 = {tint_symbol_4.position, tint_symbol_4.color, tint_symbol_4.quad_pos};
+  const VertexOutput inner_result = vs_main_inner(tint_symbol_32);
+  tint_symbol_6 wrapper_result = (tint_symbol_6)0;
+  wrapper_result.position = inner_result.position;
+  wrapper_result.color = inner_result.color;
+  wrapper_result.quad_pos = inner_result.quad_pos;
+  return wrapper_result;
+}
+
+struct tint_symbol_8 {
+  float4 color : TEXCOORD0;
+  float2 quad_pos : TEXCOORD1;
+  float4 position : SV_Position;
+};
+struct tint_symbol_9 {
+  float4 value : SV_Target0;
+};
+
+float4 fs_main_inner(VertexOutput tint_symbol) {
+  float4 color = tint_symbol.color;
+  color.a = (color.a * max((1.0f - length(tint_symbol.quad_pos)), 0.0f));
+  return color;
+}
+
+tint_symbol_9 fs_main(tint_symbol_8 tint_symbol_7) {
+  const VertexOutput tint_symbol_33 = {tint_symbol_7.position, tint_symbol_7.color, tint_symbol_7.quad_pos};
+  const float4 inner_result_1 = fs_main_inner(tint_symbol_33);
+  tint_symbol_9 wrapper_result_1 = (tint_symbol_9)0;
+  wrapper_result_1.value = inner_result_1;
+  return wrapper_result_1;
+}
+
+struct Particle {
+  float3 position;
+  float lifetime;
+  float4 color;
+  float3 velocity;
+};
+
+cbuffer cbuffer_sim_params : register(b0, space0) {
+  uint4 sim_params[2];
+};
+RWByteAddressBuffer data : register(u1, space0);
+Texture2D<float4> tint_symbol_2 : register(t2, space0);
+
+struct tint_symbol_11 {
+  uint3 GlobalInvocationID : SV_DispatchThreadID;
+};
+
+Particle tint_symbol_20(RWByteAddressBuffer buffer, uint offset) {
+  const Particle tint_symbol_34 = {asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load((offset + 12u))), asfloat(buffer.Load4((offset + 16u))), asfloat(buffer.Load3((offset + 32u)))};
+  return tint_symbol_34;
+}
+
+void tint_symbol_25(RWByteAddressBuffer buffer, uint offset, Particle value) {
+  buffer.Store3((offset + 0u), asuint(value.position));
+  buffer.Store((offset + 12u), asuint(value.lifetime));
+  buffer.Store4((offset + 16u), asuint(value.color));
+  buffer.Store3((offset + 32u), asuint(value.velocity));
+}
+
+void simulate_inner(uint3 GlobalInvocationID) {
+  rand_seed = ((asfloat(sim_params[1]).xy + float2(GlobalInvocationID.xy)) * asfloat(sim_params[1]).zw);
+  const uint idx = GlobalInvocationID.x;
+  Particle particle = tint_symbol_20(data, (48u * idx));
+  particle.velocity.z = (particle.velocity.z - (asfloat(sim_params[0].x) * 0.5f));
+  particle.position = (particle.position + (asfloat(sim_params[0].x) * particle.velocity));
+  particle.lifetime = (particle.lifetime - asfloat(sim_params[0].x));
+  particle.color.a = smoothstep(0.0f, 0.5f, particle.lifetime);
+  if ((particle.lifetime < 0.0f)) {
+    int2 coord = int2(0, 0);
+    {
+      int3 tint_tmp;
+      tint_symbol_2.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
+      int level = (tint_tmp.z - 1);
+      [loop] for(; (level > 0); level = (level - 1)) {
+        const float4 probabilites = tint_symbol_2.Load(int3(coord, level));
+        const float4 value = float4((rand()).xxxx);
+        const bool4 mask = ((value >= float4(0.0f, probabilites.xyz)) & (value < probabilites));
+        coord = (coord * 2);
+        coord.x = (coord.x + (any(mask.yw) ? 1 : 0));
+        coord.y = (coord.y + (any(mask.zw) ? 1 : 0));
+      }
+    }
+    int2 tint_tmp_1;
+    tint_symbol_2.GetDimensions(tint_tmp_1.x, tint_tmp_1.y);
+    const float2 uv = (float2(coord) / float2(tint_tmp_1));
+    particle.position = float3((((uv - 0.5f) * 3.0f) * float2(1.0f, -1.0f)), 0.0f);
+    particle.color = tint_symbol_2.Load(int3(coord, 0));
+    particle.velocity.x = ((rand() - 0.5f) * 0.100000001f);
+    particle.velocity.y = ((rand() - 0.5f) * 0.100000001f);
+    particle.velocity.z = (rand() * 0.300000012f);
+    particle.lifetime = (0.5f + (rand() * 2.0f));
+  }
+  tint_symbol_25(data, (48u * idx), particle);
+}
+
+[numthreads(64, 1, 1)]
+void simulate(tint_symbol_11 tint_symbol_10) {
+  simulate_inner(tint_symbol_10.GlobalInvocationID);
+  return;
+}
+
+cbuffer cbuffer_ubo : register(b3, space0) {
+  uint4 ubo[1];
+};
+ByteAddressBuffer buf_in : register(t4, space0);
+RWByteAddressBuffer buf_out : register(u5, space0);
+Texture2D<float4> tex_in : register(t6, space0);
+RWTexture2D<float4> tex_out : register(u7, space0);
+
+struct tint_symbol_13 {
+  uint3 coord : SV_DispatchThreadID;
+};
+
+void import_level_inner(uint3 coord) {
+  const uint offset = (coord.x + (coord.y * ubo[0].x));
+  buf_out.Store((4u * offset), asuint(tex_in.Load(int3(int2(coord.xy), 0)).w));
+}
+
+[numthreads(64, 1, 1)]
+void import_level(tint_symbol_13 tint_symbol_12) {
+  import_level_inner(tint_symbol_12.coord);
+  return;
+}
+
+struct tint_symbol_15 {
+  uint3 coord : SV_DispatchThreadID;
+};
+
+void export_level_inner(uint3 coord) {
+  int2 tint_tmp_2;
+  tex_out.GetDimensions(tint_tmp_2.x, tint_tmp_2.y);
+  if (all((coord.xy < uint2(tint_tmp_2)))) {
+    const uint dst_offset = (coord.x + (coord.y * ubo[0].x));
+    const uint src_offset = ((coord.x * 2u) + ((coord.y * 2u) * ubo[0].x));
+    const float a_1 = asfloat(buf_in.Load((4u * (src_offset + 0u))));
+    const float b = asfloat(buf_in.Load((4u * (src_offset + 1u))));
+    const float c = asfloat(buf_in.Load((4u * ((src_offset + 0u) + ubo[0].x))));
+    const float d = asfloat(buf_in.Load((4u * ((src_offset + 1u) + ubo[0].x))));
+    const float sum = dot(float4(a_1, b, c, d), float4((1.0f).xxxx));
+    buf_out.Store((4u * dst_offset), asuint((sum / 4.0f)));
+    const float4 probabilities = (float4(a_1, (a_1 + b), ((a_1 + b) + c), sum) / max(sum, 0.0001f));
+    tex_out[int2(coord.xy)] = probabilities;
+  }
+}
+
+[numthreads(64, 1, 1)]
+void export_level(tint_symbol_15 tint_symbol_14) {
+  export_level_inner(tint_symbol_14.coord);
+  return;
+}
diff --git a/test/benchmark/particles.wgsl.expected.msl b/test/benchmark/particles.wgsl.expected.msl
new file mode 100644
index 0000000..5644e7c
--- /dev/null
+++ b/test/benchmark/particles.wgsl.expected.msl
@@ -0,0 +1,175 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, int N, int M>
+inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
+struct RenderParams {
+  /* 0x0000 */ float4x4 modelViewProjectionMatrix;
+  /* 0x0040 */ packed_float3 right;
+  /* 0x004c */ int8_t tint_pad[4];
+  /* 0x0050 */ packed_float3 up;
+  /* 0x005c */ int8_t tint_pad_1[4];
+};
+struct VertexInput {
+  float3 position;
+  float4 color;
+  float2 quad_pos;
+};
+struct VertexOutput {
+  float4 position;
+  float4 color;
+  float2 quad_pos;
+};
+struct tint_symbol_2 {
+  float3 position [[attribute(0)]];
+  float4 color [[attribute(1)]];
+  float2 quad_pos [[attribute(2)]];
+};
+struct tint_symbol_3 {
+  float4 color [[user(locn0)]];
+  float2 quad_pos [[user(locn1)]];
+  float4 position [[position]];
+};
+struct tint_symbol_5 {
+  float4 color [[user(locn0)]];
+  float2 quad_pos [[user(locn1)]];
+};
+struct tint_symbol_6 {
+  float4 value [[color(0)]];
+};
+struct SimulationParams {
+  /* 0x0000 */ float deltaTime;
+  /* 0x0004 */ int8_t tint_pad_2[12];
+  /* 0x0010 */ float4 seed;
+};
+struct Particle {
+  /* 0x0000 */ packed_float3 position;
+  /* 0x000c */ float lifetime;
+  /* 0x0010 */ float4 color;
+  /* 0x0020 */ packed_float3 velocity;
+  /* 0x002c */ int8_t tint_pad_3[4];
+};
+struct Particles {
+  /* 0x0000 */ Particle particles[1];
+};
+struct UBO {
+  /* 0x0000 */ uint width;
+};
+struct Buffer {
+  /* 0x0000 */ float weights[1];
+};
+
+float rand(thread float2* const tint_symbol_9) {
+  (*(tint_symbol_9))[0] = fract((cos(dot(*(tint_symbol_9), float2(23.140779495f, 232.616897583f))) * 136.816802979f));
+  (*(tint_symbol_9))[1] = fract((cos(dot(*(tint_symbol_9), float2(54.478565216f, 345.841522217f))) * 534.764526367f));
+  return (*(tint_symbol_9))[1];
+}
+
+VertexOutput vs_main_inner(VertexInput in, const constant RenderParams* const tint_symbol_10) {
+  float3 quad_pos = (float2x3((*(tint_symbol_10)).right, (*(tint_symbol_10)).up) * in.quad_pos);
+  float3 position = (in.position + (quad_pos * 0.01f));
+  VertexOutput out = {};
+  out.position = ((*(tint_symbol_10)).modelViewProjectionMatrix * float4(position, 1.0f));
+  out.color = in.color;
+  out.quad_pos = in.quad_pos;
+  return out;
+}
+
+vertex tint_symbol_3 vs_main(const constant RenderParams* tint_symbol_11 [[buffer(0)]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  VertexInput const tint_symbol_7 = {.position=tint_symbol_1.position, .color=tint_symbol_1.color, .quad_pos=tint_symbol_1.quad_pos};
+  VertexOutput const inner_result = vs_main_inner(tint_symbol_7, tint_symbol_11);
+  tint_symbol_3 wrapper_result = {};
+  wrapper_result.position = inner_result.position;
+  wrapper_result.color = inner_result.color;
+  wrapper_result.quad_pos = inner_result.quad_pos;
+  return wrapper_result;
+}
+
+float4 fs_main_inner(VertexOutput in) {
+  float4 color = in.color;
+  color[3] = (color[3] * fmax((1.0f - length(in.quad_pos)), 0.0f));
+  return color;
+}
+
+fragment tint_symbol_6 fs_main(float4 position [[position]], tint_symbol_5 tint_symbol_4 [[stage_in]]) {
+  VertexOutput const tint_symbol_8 = {.position=position, .color=tint_symbol_4.color, .quad_pos=tint_symbol_4.quad_pos};
+  float4 const inner_result_1 = fs_main_inner(tint_symbol_8);
+  tint_symbol_6 wrapper_result_1 = {};
+  wrapper_result_1.value = inner_result_1;
+  return wrapper_result_1;
+}
+
+void simulate_inner(uint3 GlobalInvocationID, thread float2* const tint_symbol_12, const constant SimulationParams* const tint_symbol_13, device Particles* const tint_symbol_14, texture2d<float, access::sample> tint_symbol_15) {
+  *(tint_symbol_12) = ((float4((*(tint_symbol_13)).seed).xy + float2(uint3(GlobalInvocationID).xy)) * float4((*(tint_symbol_13)).seed).zw);
+  uint const idx = GlobalInvocationID[0];
+  Particle particle = (*(tint_symbol_14)).particles[idx];
+  particle.velocity[2] = (particle.velocity[2] - ((*(tint_symbol_13)).deltaTime * 0.5f));
+  particle.position = (particle.position + ((*(tint_symbol_13)).deltaTime * particle.velocity));
+  particle.lifetime = (particle.lifetime - (*(tint_symbol_13)).deltaTime);
+  particle.color[3] = smoothstep(0.0f, 0.5f, particle.lifetime);
+  if ((particle.lifetime < 0.0f)) {
+    int2 coord = int2(0, 0);
+    for(int level = as_type<int>((as_type<uint>(int(tint_symbol_15.get_num_mip_levels())) - as_type<uint>(1))); (level > 0); level = as_type<int>((as_type<uint>(level) - as_type<uint>(1)))) {
+      float4 const probabilites = tint_symbol_15.read(uint2(coord), level);
+      float4 const value = float4(rand(tint_symbol_12));
+      bool4 const mask = ((value >= float4(0.0f, float4(probabilites).xyz)) & (value < probabilites));
+      coord = as_type<int2>((as_type<uint2>(coord) * as_type<uint>(2)));
+      coord[0] = as_type<int>((as_type<uint>(coord[0]) + as_type<uint>(select(0, 1, any(bool4(mask).yw)))));
+      coord[1] = as_type<int>((as_type<uint>(coord[1]) + as_type<uint>(select(0, 1, any(bool4(mask).zw)))));
+    }
+    float2 const uv = (float2(coord) / float2(int2(tint_symbol_15.get_width(), tint_symbol_15.get_height())));
+    particle.position = float3((((uv - 0.5f) * 3.0f) * float2(1.0f, -1.0f)), 0.0f);
+    particle.color = tint_symbol_15.read(uint2(coord), 0);
+    particle.velocity[0] = ((rand(tint_symbol_12) - 0.5f) * 0.100000001f);
+    particle.velocity[1] = ((rand(tint_symbol_12) - 0.5f) * 0.100000001f);
+    particle.velocity[2] = (rand(tint_symbol_12) * 0.300000012f);
+    particle.lifetime = (0.5f + (rand(tint_symbol_12) * 2.0f));
+  }
+  (*(tint_symbol_14)).particles[idx] = particle;
+}
+
+kernel void simulate(const constant SimulationParams* tint_symbol_17 [[buffer(0)]], device Particles* tint_symbol_18 [[buffer(1)]], texture2d<float, access::sample> tint_symbol_19 [[texture(0)]], uint3 GlobalInvocationID [[thread_position_in_grid]]) {
+  thread float2 tint_symbol_16 = 0.0f;
+  simulate_inner(GlobalInvocationID, &(tint_symbol_16), tint_symbol_17, tint_symbol_18, tint_symbol_19);
+  return;
+}
+
+void import_level_inner(uint3 coord, const constant UBO* const tint_symbol_20, device Buffer* const tint_symbol_21, texture2d<float, access::sample> tint_symbol_22) {
+  uint const offset = (coord[0] + (coord[1] * (*(tint_symbol_20)).width));
+  (*(tint_symbol_21)).weights[offset] = tint_symbol_22.read(uint2(int2(uint3(coord).xy)), 0)[3];
+}
+
+kernel void import_level(const constant UBO* tint_symbol_23 [[buffer(2)]], device Buffer* tint_symbol_24 [[buffer(3)]], texture2d<float, access::sample> tint_symbol_25 [[texture(1)]], uint3 coord [[thread_position_in_grid]]) {
+  import_level_inner(coord, tint_symbol_23, tint_symbol_24, tint_symbol_25);
+  return;
+}
+
+void export_level_inner(uint3 coord, texture2d<float, access::write> tint_symbol_26, const constant UBO* const tint_symbol_27, const device Buffer* const tint_symbol_28, device Buffer* const tint_symbol_29) {
+  if (all((uint3(coord).xy < uint2(int2(tint_symbol_26.get_width(), tint_symbol_26.get_height()))))) {
+    uint const dst_offset = (coord[0] + (coord[1] * (*(tint_symbol_27)).width));
+    uint const src_offset = ((coord[0] * 2u) + ((coord[1] * 2u) * (*(tint_symbol_27)).width));
+    float const a_1 = (*(tint_symbol_28)).weights[(src_offset + 0u)];
+    float const b = (*(tint_symbol_28)).weights[(src_offset + 1u)];
+    float const c = (*(tint_symbol_28)).weights[((src_offset + 0u) + (*(tint_symbol_27)).width)];
+    float const d = (*(tint_symbol_28)).weights[((src_offset + 1u) + (*(tint_symbol_27)).width)];
+    float const sum = dot(float4(a_1, b, c, d), float4(1.0f));
+    (*(tint_symbol_29)).weights[dst_offset] = (sum / 4.0f);
+    float4 const probabilities = (float4(a_1, (a_1 + b), ((a_1 + b) + c), sum) / fmax(sum, 0.0001f));
+    tint_symbol_26.write(probabilities, uint2(int2(uint3(coord).xy)));
+  }
+}
+
+kernel void export_level(texture2d<float, access::write> tint_symbol_30 [[texture(2)]], const constant UBO* tint_symbol_31 [[buffer(2)]], const device Buffer* tint_symbol_32 [[buffer(4)]], device Buffer* tint_symbol_33 [[buffer(3)]], uint3 coord [[thread_position_in_grid]]) {
+  export_level_inner(coord, tint_symbol_30, tint_symbol_31, tint_symbol_32, tint_symbol_33);
+  return;
+}
+
diff --git a/test/benchmark/particles.wgsl.expected.spvasm b/test/benchmark/particles.wgsl.expected.spvasm
new file mode 100644
index 0000000..fe13031
--- /dev/null
+++ b/test/benchmark/particles.wgsl.expected.spvasm
@@ -0,0 +1,646 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 423
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+         %67 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %vs_main "vs_main" %position_1 %color_1 %quad_pos_1 %position_2 %color_2 %quad_pos_2 %vertex_point_size
+               OpEntryPoint Fragment %fs_main "fs_main" %position_3 %color_3 %quad_pos_3 %value_1
+               OpEntryPoint GLCompute %simulate "simulate" %GlobalInvocationID_1
+               OpEntryPoint GLCompute %import_level "import_level" %coord_1
+               OpEntryPoint GLCompute %export_level "export_level" %coord_2
+               OpExecutionMode %fs_main OriginUpperLeft
+               OpExecutionMode %simulate LocalSize 64 1 1
+               OpExecutionMode %import_level LocalSize 64 1 1
+               OpExecutionMode %export_level LocalSize 64 1 1
+               OpName %position_1 "position_1"
+               OpName %color_1 "color_1"
+               OpName %quad_pos_1 "quad_pos_1"
+               OpName %position_2 "position_2"
+               OpName %color_2 "color_2"
+               OpName %quad_pos_2 "quad_pos_2"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %position_3 "position_3"
+               OpName %color_3 "color_3"
+               OpName %quad_pos_3 "quad_pos_3"
+               OpName %value_1 "value_1"
+               OpName %GlobalInvocationID_1 "GlobalInvocationID_1"
+               OpName %coord_1 "coord_1"
+               OpName %coord_2 "coord_2"
+               OpName %rand_seed "rand_seed"
+               OpName %RenderParams "RenderParams"
+               OpMemberName %RenderParams 0 "modelViewProjectionMatrix"
+               OpMemberName %RenderParams 1 "right"
+               OpMemberName %RenderParams 2 "up"
+               OpName %render_params "render_params"
+               OpName %SimulationParams "SimulationParams"
+               OpMemberName %SimulationParams 0 "deltaTime"
+               OpMemberName %SimulationParams 1 "seed"
+               OpName %sim_params "sim_params"
+               OpName %Particles "Particles"
+               OpMemberName %Particles 0 "particles"
+               OpName %Particle "Particle"
+               OpMemberName %Particle 0 "position"
+               OpMemberName %Particle 1 "lifetime"
+               OpMemberName %Particle 2 "color"
+               OpMemberName %Particle 3 "velocity"
+               OpName %data "data"
+               OpName %texture "texture"
+               OpName %UBO "UBO"
+               OpMemberName %UBO 0 "width"
+               OpName %ubo "ubo"
+               OpName %Buffer "Buffer"
+               OpMemberName %Buffer 0 "weights"
+               OpName %buf_in "buf_in"
+               OpName %buf_out "buf_out"
+               OpName %tex_in "tex_in"
+               OpName %tex_out "tex_out"
+               OpName %rand "rand"
+               OpName %VertexOutput "VertexOutput"
+               OpMemberName %VertexOutput 0 "position"
+               OpMemberName %VertexOutput 1 "color"
+               OpMemberName %VertexOutput 2 "quad_pos"
+               OpName %VertexInput "VertexInput"
+               OpMemberName %VertexInput 0 "position"
+               OpMemberName %VertexInput 1 "color"
+               OpMemberName %VertexInput 2 "quad_pos"
+               OpName %vs_main_inner "vs_main_inner"
+               OpName %in "in"
+               OpName %quad_pos "quad_pos"
+               OpName %position "position"
+               OpName %out "out"
+               OpName %vs_main "vs_main"
+               OpName %fs_main_inner "fs_main_inner"
+               OpName %in_0 "in"
+               OpName %color "color"
+               OpName %fs_main "fs_main"
+               OpName %simulate_inner "simulate_inner"
+               OpName %GlobalInvocationID "GlobalInvocationID"
+               OpName %particle "particle"
+               OpName %coord "coord"
+               OpName %level "level"
+               OpName %simulate "simulate"
+               OpName %import_level_inner "import_level_inner"
+               OpName %coord_0 "coord"
+               OpName %import_level "import_level"
+               OpName %export_level_inner "export_level_inner"
+               OpName %coord_3 "coord"
+               OpName %export_level "export_level"
+               OpDecorate %position_1 Location 0
+               OpDecorate %color_1 Location 1
+               OpDecorate %quad_pos_1 Location 2
+               OpDecorate %position_2 BuiltIn Position
+               OpDecorate %color_2 Location 0
+               OpDecorate %quad_pos_2 Location 1
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpDecorate %position_3 BuiltIn FragCoord
+               OpDecorate %color_3 Location 0
+               OpDecorate %quad_pos_3 Location 1
+               OpDecorate %value_1 Location 0
+               OpDecorate %GlobalInvocationID_1 BuiltIn GlobalInvocationId
+               OpDecorate %coord_1 BuiltIn GlobalInvocationId
+               OpDecorate %coord_2 BuiltIn GlobalInvocationId
+               OpDecorate %RenderParams Block
+               OpMemberDecorate %RenderParams 0 Offset 0
+               OpMemberDecorate %RenderParams 0 ColMajor
+               OpMemberDecorate %RenderParams 0 MatrixStride 16
+               OpMemberDecorate %RenderParams 1 Offset 64
+               OpMemberDecorate %RenderParams 2 Offset 80
+               OpDecorate %render_params NonWritable
+               OpDecorate %render_params Binding 0
+               OpDecorate %render_params DescriptorSet 0
+               OpDecorate %SimulationParams Block
+               OpMemberDecorate %SimulationParams 0 Offset 0
+               OpMemberDecorate %SimulationParams 1 Offset 16
+               OpDecorate %sim_params NonWritable
+               OpDecorate %sim_params Binding 0
+               OpDecorate %sim_params DescriptorSet 0
+               OpDecorate %Particles Block
+               OpMemberDecorate %Particles 0 Offset 0
+               OpMemberDecorate %Particle 0 Offset 0
+               OpMemberDecorate %Particle 1 Offset 12
+               OpMemberDecorate %Particle 2 Offset 16
+               OpMemberDecorate %Particle 3 Offset 32
+               OpDecorate %_runtimearr_Particle ArrayStride 48
+               OpDecorate %data Binding 1
+               OpDecorate %data DescriptorSet 0
+               OpDecorate %texture Binding 2
+               OpDecorate %texture DescriptorSet 0
+               OpDecorate %UBO Block
+               OpMemberDecorate %UBO 0 Offset 0
+               OpDecorate %ubo NonWritable
+               OpDecorate %ubo Binding 3
+               OpDecorate %ubo DescriptorSet 0
+               OpDecorate %Buffer Block
+               OpMemberDecorate %Buffer 0 Offset 0
+               OpDecorate %_runtimearr_float ArrayStride 4
+               OpDecorate %buf_in NonWritable
+               OpDecorate %buf_in Binding 4
+               OpDecorate %buf_in DescriptorSet 0
+               OpDecorate %buf_out Binding 5
+               OpDecorate %buf_out DescriptorSet 0
+               OpDecorate %tex_in Binding 6
+               OpDecorate %tex_in DescriptorSet 0
+               OpDecorate %tex_out NonReadable
+               OpDecorate %tex_out Binding 7
+               OpDecorate %tex_out DescriptorSet 0
+               OpMemberDecorate %VertexOutput 0 Offset 0
+               OpMemberDecorate %VertexOutput 1 Offset 16
+               OpMemberDecorate %VertexOutput 2 Offset 32
+               OpMemberDecorate %VertexInput 0 Offset 0
+               OpMemberDecorate %VertexInput 1 Offset 16
+               OpMemberDecorate %VertexInput 2 Offset 32
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%_ptr_Input_v3float = OpTypePointer Input %v3float
+ %position_1 = OpVariable %_ptr_Input_v3float Input
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+    %color_1 = OpVariable %_ptr_Input_v4float Input
+    %v2float = OpTypeVector %float 2
+%_ptr_Input_v2float = OpTypePointer Input %v2float
+ %quad_pos_1 = OpVariable %_ptr_Input_v2float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+         %13 = OpConstantNull %v4float
+ %position_2 = OpVariable %_ptr_Output_v4float Output %13
+    %color_2 = OpVariable %_ptr_Output_v4float Output %13
+%_ptr_Output_v2float = OpTypePointer Output %v2float
+         %17 = OpConstantNull %v2float
+ %quad_pos_2 = OpVariable %_ptr_Output_v2float Output %17
+%_ptr_Output_float = OpTypePointer Output %float
+         %20 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %20
+ %position_3 = OpVariable %_ptr_Input_v4float Input
+    %color_3 = OpVariable %_ptr_Input_v4float Input
+ %quad_pos_3 = OpVariable %_ptr_Input_v2float Input
+    %value_1 = OpVariable %_ptr_Output_v4float Output %13
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%GlobalInvocationID_1 = OpVariable %_ptr_Input_v3uint Input
+    %coord_1 = OpVariable %_ptr_Input_v3uint Input
+    %coord_2 = OpVariable %_ptr_Input_v3uint Input
+%_ptr_Private_v2float = OpTypePointer Private %v2float
+  %rand_seed = OpVariable %_ptr_Private_v2float Private %17
+%mat4v4float = OpTypeMatrix %v4float 4
+%RenderParams = OpTypeStruct %mat4v4float %v3float %v3float
+%_ptr_Uniform_RenderParams = OpTypePointer Uniform %RenderParams
+%render_params = OpVariable %_ptr_Uniform_RenderParams Uniform
+%SimulationParams = OpTypeStruct %float %v4float
+%_ptr_Uniform_SimulationParams = OpTypePointer Uniform %SimulationParams
+ %sim_params = OpVariable %_ptr_Uniform_SimulationParams Uniform
+   %Particle = OpTypeStruct %v3float %float %v4float %v3float
+%_runtimearr_Particle = OpTypeRuntimeArray %Particle
+  %Particles = OpTypeStruct %_runtimearr_Particle
+%_ptr_StorageBuffer_Particles = OpTypePointer StorageBuffer %Particles
+       %data = OpVariable %_ptr_StorageBuffer_Particles StorageBuffer
+         %47 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_47 = OpTypePointer UniformConstant %47
+    %texture = OpVariable %_ptr_UniformConstant_47 UniformConstant
+        %UBO = OpTypeStruct %uint
+%_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
+        %ubo = OpVariable %_ptr_Uniform_UBO Uniform
+%_runtimearr_float = OpTypeRuntimeArray %float
+     %Buffer = OpTypeStruct %_runtimearr_float
+%_ptr_StorageBuffer_Buffer = OpTypePointer StorageBuffer %Buffer
+     %buf_in = OpVariable %_ptr_StorageBuffer_Buffer StorageBuffer
+    %buf_out = OpVariable %_ptr_StorageBuffer_Buffer StorageBuffer
+     %tex_in = OpVariable %_ptr_UniformConstant_47 UniformConstant
+         %59 = OpTypeImage %float 2D 0 0 0 2 Rgba8
+%_ptr_UniformConstant_59 = OpTypePointer UniformConstant %59
+    %tex_out = OpVariable %_ptr_UniformConstant_59 UniformConstant
+         %60 = OpTypeFunction %float
+     %uint_0 = OpConstant %uint 0
+%_ptr_Private_float = OpTypePointer Private %float
+%float_23_1407795 = OpConstant %float 23.1407795
+%float_232_616898 = OpConstant %float 232.616898
+         %73 = OpConstantComposite %v2float %float_23_1407795 %float_232_616898
+%float_136_816803 = OpConstant %float 136.816803
+     %uint_1 = OpConstant %uint 1
+%float_54_4785652 = OpConstant %float 54.4785652
+%float_345_841522 = OpConstant %float 345.841522
+         %84 = OpConstantComposite %v2float %float_54_4785652 %float_345_841522
+%float_534_764526 = OpConstant %float 534.764526
+%VertexOutput = OpTypeStruct %v4float %v4float %v2float
+%VertexInput = OpTypeStruct %v3float %v4float %v2float
+         %89 = OpTypeFunction %VertexOutput %VertexInput
+%mat2v3float = OpTypeMatrix %v3float 2
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+     %uint_2 = OpConstant %uint 2
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+        %107 = OpConstantNull %v3float
+%float_0_00999999978 = OpConstant %float 0.00999999978
+%_ptr_Function_VertexOutput = OpTypePointer Function %VertexOutput
+        %116 = OpConstantNull %VertexOutput
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+%_ptr_Uniform_mat4v4float = OpTypePointer Uniform %mat4v4float
+    %float_1 = OpConstant %float 1
+%_ptr_Function_v2float = OpTypePointer Function %v2float
+       %void = OpTypeVoid
+        %135 = OpTypeFunction %void
+        %147 = OpTypeFunction %v4float %VertexOutput
+     %uint_3 = OpConstant %uint 3
+%_ptr_Function_float = OpTypePointer Function %float
+    %float_0 = OpConstant %float 0
+        %172 = OpTypeFunction %void %v3uint
+%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
+     %v2uint = OpTypeVector %uint 2
+%_ptr_StorageBuffer_Particle = OpTypePointer StorageBuffer %Particle
+%_ptr_Function_Particle = OpTypePointer Function %Particle
+        %194 = OpConstantNull %Particle
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+  %float_0_5 = OpConstant %float 0.5
+       %bool = OpTypeBool
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+      %int_0 = OpConstant %int 0
+        %232 = OpConstantComposite %v2int %int_0 %int_0
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+        %235 = OpConstantNull %v2int
+      %int_1 = OpConstant %int 1
+%_ptr_Function_int = OpTypePointer Function %int
+        %242 = OpConstantNull %int
+     %v4bool = OpTypeVector %bool 4
+      %int_2 = OpConstant %int 2
+     %v2bool = OpTypeVector %bool 2
+    %float_3 = OpConstant %float 3
+   %float_n1 = OpConstant %float -1
+        %302 = OpConstantComposite %v2float %float_1 %float_n1
+%float_0_100000001 = OpConstant %float 0.100000001
+%float_0_300000012 = OpConstant %float 0.300000012
+    %float_2 = OpConstant %float 2
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+        %402 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+    %float_4 = OpConstant %float 4
+%float_9_99999975en05 = OpConstant %float 9.99999975e-05
+       %rand = OpFunction %float None %60
+         %62 = OpLabel
+         %65 = OpAccessChain %_ptr_Private_float %rand_seed %uint_0
+         %70 = OpLoad %v2float %rand_seed
+         %69 = OpDot %float %70 %73
+         %68 = OpExtInst %float %67 Cos %69
+         %75 = OpFMul %float %68 %float_136_816803
+         %66 = OpExtInst %float %67 Fract %75
+               OpStore %65 %66
+         %77 = OpAccessChain %_ptr_Private_float %rand_seed %uint_1
+         %81 = OpLoad %v2float %rand_seed
+         %80 = OpDot %float %81 %84
+         %79 = OpExtInst %float %67 Cos %80
+         %86 = OpFMul %float %79 %float_534_764526
+         %78 = OpExtInst %float %67 Fract %86
+               OpStore %77 %78
+         %87 = OpAccessChain %_ptr_Private_float %rand_seed %uint_1
+         %88 = OpLoad %float %87
+               OpReturnValue %88
+               OpFunctionEnd
+%vs_main_inner = OpFunction %VertexOutput None %89
+         %in = OpFunctionParameter %VertexInput
+         %94 = OpLabel
+   %quad_pos = OpVariable %_ptr_Function_v3float Function %107
+   %position = OpVariable %_ptr_Function_v3float Function %107
+        %out = OpVariable %_ptr_Function_VertexOutput Function %116
+         %97 = OpAccessChain %_ptr_Uniform_v3float %render_params %uint_1
+         %98 = OpLoad %v3float %97
+        %100 = OpAccessChain %_ptr_Uniform_v3float %render_params %uint_2
+        %101 = OpLoad %v3float %100
+        %102 = OpCompositeConstruct %mat2v3float %98 %101
+        %103 = OpCompositeExtract %v2float %in 2
+        %104 = OpMatrixTimesVector %v3float %102 %103
+               OpStore %quad_pos %104
+        %108 = OpCompositeExtract %v3float %in 0
+        %109 = OpLoad %v3float %quad_pos
+        %111 = OpVectorTimesScalar %v3float %109 %float_0_00999999978
+        %112 = OpFAdd %v3float %108 %111
+               OpStore %position %112
+        %118 = OpAccessChain %_ptr_Function_v4float %out %uint_0
+        %120 = OpAccessChain %_ptr_Uniform_mat4v4float %render_params %uint_0
+        %121 = OpLoad %mat4v4float %120
+        %122 = OpLoad %v3float %position
+        %123 = OpCompositeExtract %float %122 0
+        %124 = OpCompositeExtract %float %122 1
+        %125 = OpCompositeExtract %float %122 2
+        %127 = OpCompositeConstruct %v4float %123 %124 %125 %float_1
+        %128 = OpMatrixTimesVector %v4float %121 %127
+               OpStore %118 %128
+        %129 = OpAccessChain %_ptr_Function_v4float %out %uint_1
+        %130 = OpCompositeExtract %v4float %in 1
+               OpStore %129 %130
+        %132 = OpAccessChain %_ptr_Function_v2float %out %uint_2
+        %133 = OpCompositeExtract %v2float %in 2
+               OpStore %132 %133
+        %134 = OpLoad %VertexOutput %out
+               OpReturnValue %134
+               OpFunctionEnd
+    %vs_main = OpFunction %void None %135
+        %138 = OpLabel
+        %140 = OpLoad %v3float %position_1
+        %141 = OpLoad %v4float %color_1
+        %142 = OpLoad %v2float %quad_pos_1
+        %143 = OpCompositeConstruct %VertexInput %140 %141 %142
+        %139 = OpFunctionCall %VertexOutput %vs_main_inner %143
+        %144 = OpCompositeExtract %v4float %139 0
+               OpStore %position_2 %144
+        %145 = OpCompositeExtract %v4float %139 1
+               OpStore %color_2 %145
+        %146 = OpCompositeExtract %v2float %139 2
+               OpStore %quad_pos_2 %146
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
+%fs_main_inner = OpFunction %v4float None %147
+       %in_0 = OpFunctionParameter %VertexOutput
+        %150 = OpLabel
+      %color = OpVariable %_ptr_Function_v4float Function %13
+        %151 = OpCompositeExtract %v4float %in_0 1
+               OpStore %color %151
+        %155 = OpAccessChain %_ptr_Function_float %color %uint_3
+        %156 = OpAccessChain %_ptr_Function_float %color %uint_3
+        %157 = OpLoad %float %156
+        %160 = OpCompositeExtract %v2float %in_0 2
+        %159 = OpExtInst %float %67 Length %160
+        %161 = OpFSub %float %float_1 %159
+        %158 = OpExtInst %float %67 NMax %161 %float_0
+        %163 = OpFMul %float %157 %158
+               OpStore %155 %163
+        %164 = OpLoad %v4float %color
+               OpReturnValue %164
+               OpFunctionEnd
+    %fs_main = OpFunction %void None %135
+        %166 = OpLabel
+        %168 = OpLoad %v4float %position_3
+        %169 = OpLoad %v4float %color_3
+        %170 = OpLoad %v2float %quad_pos_3
+        %171 = OpCompositeConstruct %VertexOutput %168 %169 %170
+        %167 = OpFunctionCall %v4float %fs_main_inner %171
+               OpStore %value_1 %167
+               OpReturn
+               OpFunctionEnd
+%simulate_inner = OpFunction %void None %172
+%GlobalInvocationID = OpFunctionParameter %v3uint
+        %175 = OpLabel
+   %particle = OpVariable %_ptr_Function_Particle Function %194
+      %coord = OpVariable %_ptr_Function_v2int Function %235
+      %level = OpVariable %_ptr_Function_int Function %242
+        %270 = OpVariable %_ptr_Function_v2int Function %235
+        %297 = OpVariable %_ptr_Function_v2float Function %17
+        %177 = OpAccessChain %_ptr_Uniform_v4float %sim_params %uint_1
+        %178 = OpLoad %v4float %177
+        %179 = OpVectorShuffle %v2float %178 %178 0 1
+        %182 = OpVectorShuffle %v2uint %GlobalInvocationID %GlobalInvocationID 0 1
+        %180 = OpConvertUToF %v2float %182
+        %183 = OpFAdd %v2float %179 %180
+        %184 = OpAccessChain %_ptr_Uniform_v4float %sim_params %uint_1
+        %185 = OpLoad %v4float %184
+        %186 = OpVectorShuffle %v2float %185 %185 2 3
+        %187 = OpFMul %v2float %183 %186
+               OpStore %rand_seed %187
+        %188 = OpCompositeExtract %uint %GlobalInvocationID 0
+        %190 = OpAccessChain %_ptr_StorageBuffer_Particle %data %uint_0 %188
+        %191 = OpLoad %Particle %190
+               OpStore %particle %191
+        %195 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_2
+        %196 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_2
+        %197 = OpLoad %float %196
+        %199 = OpAccessChain %_ptr_Uniform_float %sim_params %uint_0
+        %200 = OpLoad %float %199
+        %202 = OpFMul %float %200 %float_0_5
+        %203 = OpFSub %float %197 %202
+               OpStore %195 %203
+        %204 = OpAccessChain %_ptr_Function_v3float %particle %uint_0
+        %205 = OpAccessChain %_ptr_Function_v3float %particle %uint_0
+        %206 = OpLoad %v3float %205
+        %207 = OpAccessChain %_ptr_Uniform_float %sim_params %uint_0
+        %208 = OpLoad %float %207
+        %209 = OpAccessChain %_ptr_Function_v3float %particle %uint_3
+        %210 = OpLoad %v3float %209
+        %211 = OpVectorTimesScalar %v3float %210 %208
+        %212 = OpFAdd %v3float %206 %211
+               OpStore %204 %212
+        %213 = OpAccessChain %_ptr_Function_float %particle %uint_1
+        %214 = OpAccessChain %_ptr_Function_float %particle %uint_1
+        %215 = OpLoad %float %214
+        %216 = OpAccessChain %_ptr_Uniform_float %sim_params %uint_0
+        %217 = OpLoad %float %216
+        %218 = OpFSub %float %215 %217
+               OpStore %213 %218
+        %219 = OpAccessChain %_ptr_Function_float %particle %uint_2 %uint_3
+        %221 = OpAccessChain %_ptr_Function_float %particle %uint_1
+        %222 = OpLoad %float %221
+        %220 = OpExtInst %float %67 SmoothStep %float_0 %float_0_5 %222
+               OpStore %219 %220
+        %223 = OpAccessChain %_ptr_Function_float %particle %uint_1
+        %224 = OpLoad %float %223
+        %225 = OpFOrdLessThan %bool %224 %float_0
+               OpSelectionMerge %227 None
+               OpBranchConditional %225 %228 %227
+        %228 = OpLabel
+               OpStore %coord %232
+        %237 = OpLoad %47 %texture
+        %236 = OpImageQueryLevels %int %237
+        %239 = OpISub %int %236 %int_1
+               OpStore %level %239
+               OpBranch %243
+        %243 = OpLabel
+               OpLoopMerge %244 %245 None
+               OpBranch %246
+        %246 = OpLabel
+        %248 = OpLoad %int %level
+        %249 = OpSGreaterThan %bool %248 %int_0
+        %247 = OpLogicalNot %bool %249
+               OpSelectionMerge %250 None
+               OpBranchConditional %247 %251 %250
+        %251 = OpLabel
+               OpBranch %244
+        %250 = OpLabel
+        %253 = OpLoad %47 %texture
+        %254 = OpLoad %v2int %coord
+        %255 = OpLoad %int %level
+        %252 = OpImageFetch %v4float %253 %254 Lod %255
+        %256 = OpFunctionCall %float %rand
+        %257 = OpCompositeConstruct %v4float %256 %256 %256 %256
+        %258 = OpVectorShuffle %v3float %252 %252 0 1 2
+        %259 = OpCompositeExtract %float %258 0
+        %260 = OpCompositeExtract %float %258 1
+        %261 = OpCompositeExtract %float %258 2
+        %262 = OpCompositeConstruct %v4float %float_0 %259 %260 %261
+        %263 = OpFOrdGreaterThanEqual %v4bool %257 %262
+        %265 = OpFOrdLessThan %v4bool %257 %252
+        %266 = OpLogicalAnd %v4bool %263 %265
+        %267 = OpLoad %v2int %coord
+        %271 = OpCompositeConstruct %v2int %int_2 %int_2
+        %269 = OpIMul %v2int %267 %271
+               OpStore %coord %269
+        %272 = OpAccessChain %_ptr_Function_int %coord %uint_0
+        %273 = OpAccessChain %_ptr_Function_int %coord %uint_0
+        %274 = OpLoad %int %273
+        %278 = OpVectorShuffle %v2bool %266 %266 1 3
+        %276 = OpAny %bool %278
+        %275 = OpSelect %int %276 %int_1 %int_0
+        %279 = OpIAdd %int %274 %275
+               OpStore %272 %279
+        %280 = OpAccessChain %_ptr_Function_int %coord %uint_1
+        %281 = OpAccessChain %_ptr_Function_int %coord %uint_1
+        %282 = OpLoad %int %281
+        %285 = OpVectorShuffle %v2bool %266 %266 2 3
+        %284 = OpAny %bool %285
+        %283 = OpSelect %int %284 %int_1 %int_0
+        %286 = OpIAdd %int %282 %283
+               OpStore %280 %286
+               OpBranch %245
+        %245 = OpLabel
+        %287 = OpLoad %int %level
+        %288 = OpISub %int %287 %int_1
+               OpStore %level %288
+               OpBranch %243
+        %244 = OpLabel
+        %290 = OpLoad %v2int %coord
+        %289 = OpConvertSToF %v2float %290
+        %293 = OpLoad %47 %texture
+        %292 = OpImageQuerySizeLod %v2int %293 %int_0
+        %291 = OpConvertSToF %v2float %292
+        %294 = OpFDiv %v2float %289 %291
+        %295 = OpAccessChain %_ptr_Function_v3float %particle %uint_0
+        %298 = OpCompositeConstruct %v2float %float_0_5 %float_0_5
+        %296 = OpFSub %v2float %294 %298
+        %300 = OpVectorTimesScalar %v2float %296 %float_3
+        %303 = OpFMul %v2float %300 %302
+        %304 = OpCompositeExtract %float %303 0
+        %305 = OpCompositeExtract %float %303 1
+        %306 = OpCompositeConstruct %v3float %304 %305 %float_0
+               OpStore %295 %306
+        %307 = OpAccessChain %_ptr_Function_v4float %particle %uint_2
+        %309 = OpLoad %47 %texture
+        %310 = OpLoad %v2int %coord
+        %308 = OpImageFetch %v4float %309 %310 Lod %int_0
+               OpStore %307 %308
+        %311 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_0
+        %312 = OpFunctionCall %float %rand
+        %313 = OpFSub %float %312 %float_0_5
+        %315 = OpFMul %float %313 %float_0_100000001
+               OpStore %311 %315
+        %316 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_1
+        %317 = OpFunctionCall %float %rand
+        %318 = OpFSub %float %317 %float_0_5
+        %319 = OpFMul %float %318 %float_0_100000001
+               OpStore %316 %319
+        %320 = OpAccessChain %_ptr_Function_float %particle %uint_3 %uint_2
+        %321 = OpFunctionCall %float %rand
+        %323 = OpFMul %float %321 %float_0_300000012
+               OpStore %320 %323
+        %324 = OpAccessChain %_ptr_Function_float %particle %uint_1
+        %325 = OpFunctionCall %float %rand
+        %327 = OpFMul %float %325 %float_2
+        %328 = OpFAdd %float %float_0_5 %327
+               OpStore %324 %328
+               OpBranch %227
+        %227 = OpLabel
+        %329 = OpAccessChain %_ptr_StorageBuffer_Particle %data %uint_0 %188
+        %330 = OpLoad %Particle %particle
+               OpStore %329 %330
+               OpReturn
+               OpFunctionEnd
+   %simulate = OpFunction %void None %135
+        %332 = OpLabel
+        %334 = OpLoad %v3uint %GlobalInvocationID_1
+        %333 = OpFunctionCall %void %simulate_inner %334
+               OpReturn
+               OpFunctionEnd
+%import_level_inner = OpFunction %void None %172
+    %coord_0 = OpFunctionParameter %v3uint
+        %337 = OpLabel
+        %339 = OpCompositeExtract %uint %coord_0 0
+        %340 = OpCompositeExtract %uint %coord_0 1
+        %342 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+        %343 = OpLoad %uint %342
+        %344 = OpIMul %uint %340 %343
+        %345 = OpIAdd %uint %339 %344
+        %347 = OpAccessChain %_ptr_StorageBuffer_float %buf_out %uint_0 %345
+        %349 = OpLoad %47 %tex_in
+        %351 = OpVectorShuffle %v2uint %coord_0 %coord_0 0 1
+        %350 = OpBitcast %v2int %351
+        %348 = OpImageFetch %v4float %349 %350 Lod %int_0
+        %352 = OpCompositeExtract %float %348 3
+               OpStore %347 %352
+               OpReturn
+               OpFunctionEnd
+%import_level = OpFunction %void None %135
+        %354 = OpLabel
+        %356 = OpLoad %v3uint %coord_1
+        %355 = OpFunctionCall %void %import_level_inner %356
+               OpReturn
+               OpFunctionEnd
+%export_level_inner = OpFunction %void None %172
+    %coord_3 = OpFunctionParameter %v3uint
+        %359 = OpLabel
+        %413 = OpVariable %_ptr_Function_v4float Function %13
+        %361 = OpVectorShuffle %v2uint %coord_3 %coord_3 0 1
+        %364 = OpLoad %59 %tex_out
+        %363 = OpImageQuerySize %v2int %364
+        %362 = OpBitcast %v2uint %363
+        %365 = OpULessThan %v2bool %361 %362
+        %360 = OpAll %bool %365
+               OpSelectionMerge %366 None
+               OpBranchConditional %360 %367 %366
+        %367 = OpLabel
+        %368 = OpCompositeExtract %uint %coord_3 0
+        %369 = OpCompositeExtract %uint %coord_3 1
+        %370 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+        %371 = OpLoad %uint %370
+        %372 = OpIMul %uint %369 %371
+        %373 = OpIAdd %uint %368 %372
+        %374 = OpCompositeExtract %uint %coord_3 0
+        %375 = OpIMul %uint %374 %uint_2
+        %376 = OpCompositeExtract %uint %coord_3 1
+        %377 = OpIMul %uint %376 %uint_2
+        %378 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+        %379 = OpLoad %uint %378
+        %380 = OpIMul %uint %377 %379
+        %381 = OpIAdd %uint %375 %380
+        %382 = OpIAdd %uint %381 %uint_0
+        %383 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %382
+        %384 = OpLoad %float %383
+        %385 = OpIAdd %uint %381 %uint_1
+        %386 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %385
+        %387 = OpLoad %float %386
+        %388 = OpIAdd %uint %381 %uint_0
+        %389 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+        %390 = OpLoad %uint %389
+        %391 = OpIAdd %uint %388 %390
+        %392 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %391
+        %393 = OpLoad %float %392
+        %394 = OpIAdd %uint %381 %uint_1
+        %395 = OpAccessChain %_ptr_Uniform_uint %ubo %uint_0
+        %396 = OpLoad %uint %395
+        %397 = OpIAdd %uint %394 %396
+        %398 = OpAccessChain %_ptr_StorageBuffer_float %buf_in %uint_0 %397
+        %399 = OpLoad %float %398
+        %401 = OpCompositeConstruct %v4float %384 %387 %393 %399
+        %400 = OpDot %float %401 %402
+        %403 = OpAccessChain %_ptr_StorageBuffer_float %buf_out %uint_0 %373
+        %405 = OpFDiv %float %400 %float_4
+               OpStore %403 %405
+        %406 = OpFAdd %float %384 %387
+        %407 = OpFAdd %float %384 %387
+        %408 = OpFAdd %float %407 %393
+        %409 = OpCompositeConstruct %v4float %384 %406 %408 %400
+        %410 = OpExtInst %float %67 NMax %400 %float_9_99999975en05
+        %414 = OpCompositeConstruct %v4float %410 %410 %410 %410
+        %412 = OpFDiv %v4float %409 %414
+        %416 = OpLoad %59 %tex_out
+        %418 = OpVectorShuffle %v2uint %coord_3 %coord_3 0 1
+        %417 = OpBitcast %v2int %418
+               OpImageWrite %416 %417 %412
+               OpBranch %366
+        %366 = OpLabel
+               OpReturn
+               OpFunctionEnd
+%export_level = OpFunction %void None %135
+        %420 = OpLabel
+        %422 = OpLoad %v3uint %coord_2
+        %421 = OpFunctionCall %void %export_level_inner %422
+               OpReturn
+               OpFunctionEnd
diff --git a/test/benchmark/particles.wgsl.expected.wgsl b/test/benchmark/particles.wgsl.expected.wgsl
new file mode 100644
index 0000000..856f496
--- /dev/null
+++ b/test/benchmark/particles.wgsl.expected.wgsl
@@ -0,0 +1,144 @@
+var<private> rand_seed : vec2<f32>;
+
+fn rand() -> f32 {
+  rand_seed.x = fract((cos(dot(rand_seed, vec2<f32>(23.140779495, 232.616897583))) * 136.816802979));
+  rand_seed.y = fract((cos(dot(rand_seed, vec2<f32>(54.478565216, 345.841522217))) * 534.764526367));
+  return rand_seed.y;
+}
+
+struct RenderParams {
+  modelViewProjectionMatrix : mat4x4<f32>;
+  right : vec3<f32>;
+  up : vec3<f32>;
+};
+
+[[binding(0), group(0)]] var<uniform> render_params : RenderParams;
+
+struct VertexInput {
+  [[location(0)]]
+  position : vec3<f32>;
+  [[location(1)]]
+  color : vec4<f32>;
+  [[location(2)]]
+  quad_pos : vec2<f32>;
+};
+
+struct VertexOutput {
+  [[builtin(position)]]
+  position : vec4<f32>;
+  [[location(0)]]
+  color : vec4<f32>;
+  [[location(1)]]
+  quad_pos : vec2<f32>;
+};
+
+[[stage(vertex)]]
+fn vs_main(in : VertexInput) -> VertexOutput {
+  var quad_pos = (mat2x3<f32>(render_params.right, render_params.up) * in.quad_pos);
+  var position = (in.position + (quad_pos * 0.01));
+  var out : VertexOutput;
+  out.position = (render_params.modelViewProjectionMatrix * vec4<f32>(position, 1.0));
+  out.color = in.color;
+  out.quad_pos = in.quad_pos;
+  return out;
+}
+
+[[stage(fragment)]]
+fn fs_main(in : VertexOutput) -> [[location(0)]] vec4<f32> {
+  var color = in.color;
+  color.a = (color.a * max((1.0 - length(in.quad_pos)), 0.0));
+  return color;
+}
+
+struct SimulationParams {
+  deltaTime : f32;
+  seed : vec4<f32>;
+};
+
+struct Particle {
+  position : vec3<f32>;
+  lifetime : f32;
+  color : vec4<f32>;
+  velocity : vec3<f32>;
+};
+
+struct Particles {
+  particles : array<Particle>;
+};
+
+[[binding(0), group(0)]] var<uniform> sim_params : SimulationParams;
+
+[[binding(1), group(0)]] var<storage, read_write> data : Particles;
+
+[[binding(2), group(0)]] var texture : texture_2d<f32>;
+
+[[stage(compute), workgroup_size(64)]]
+fn simulate([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+  rand_seed = ((sim_params.seed.xy + vec2<f32>(GlobalInvocationID.xy)) * sim_params.seed.zw);
+  let idx = GlobalInvocationID.x;
+  var particle = data.particles[idx];
+  particle.velocity.z = (particle.velocity.z - (sim_params.deltaTime * 0.5));
+  particle.position = (particle.position + (sim_params.deltaTime * particle.velocity));
+  particle.lifetime = (particle.lifetime - sim_params.deltaTime);
+  particle.color.a = smoothStep(0.0, 0.5, particle.lifetime);
+  if ((particle.lifetime < 0.0)) {
+    var coord = vec2<i32>(0, 0);
+    for(var level = (textureNumLevels(texture) - 1); (level > 0); level = (level - 1)) {
+      let probabilites = textureLoad(texture, coord, level);
+      let value = vec4<f32>(rand());
+      let mask = ((value >= vec4<f32>(0.0, probabilites.xyz)) & (value < probabilites));
+      coord = (coord * 2);
+      coord.x = (coord.x + select(0, 1, any(mask.yw)));
+      coord.y = (coord.y + select(0, 1, any(mask.zw)));
+    }
+    let uv = (vec2<f32>(coord) / vec2<f32>(textureDimensions(texture)));
+    particle.position = vec3<f32>((((uv - 0.5) * 3.0) * vec2<f32>(1.0, -1.0)), 0.0);
+    particle.color = textureLoad(texture, coord, 0);
+    particle.velocity.x = ((rand() - 0.5) * 0.100000001);
+    particle.velocity.y = ((rand() - 0.5) * 0.100000001);
+    particle.velocity.z = (rand() * 0.300000012);
+    particle.lifetime = (0.5 + (rand() * 2.0));
+  }
+  data.particles[idx] = particle;
+}
+
+struct UBO {
+  width : u32;
+};
+
+struct Buffer {
+  weights : array<f32>;
+};
+
+[[binding(3), group(0)]] var<uniform> ubo : UBO;
+
+[[binding(4), group(0)]] var<storage, read> buf_in : Buffer;
+
+[[binding(5), group(0)]] var<storage, read_write> buf_out : Buffer;
+
+[[binding(6), group(0)]] var tex_in : texture_2d<f32>;
+
+[[binding(7), group(0)]] var tex_out : texture_storage_2d<rgba8unorm, write>;
+
+[[stage(compute), workgroup_size(64)]]
+fn import_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+  _ = &(buf_in);
+  let offset = (coord.x + (coord.y * ubo.width));
+  buf_out.weights[offset] = textureLoad(tex_in, vec2<i32>(coord.xy), 0).w;
+}
+
+[[stage(compute), workgroup_size(64)]]
+fn export_level([[builtin(global_invocation_id)]] coord : vec3<u32>) {
+  if (all((coord.xy < vec2<u32>(textureDimensions(tex_out))))) {
+    let dst_offset = (coord.x + (coord.y * ubo.width));
+    let src_offset = ((coord.x * 2u) + ((coord.y * 2u) * ubo.width));
+    let a = buf_in.weights[(src_offset + 0u)];
+    let b = buf_in.weights[(src_offset + 1u)];
+    let c = buf_in.weights[((src_offset + 0u) + ubo.width)];
+    let d = buf_in.weights[((src_offset + 1u) + ubo.width)];
+    let sum = dot(vec4<f32>(a, b, c, d), vec4<f32>(1.0));
+    buf_out.weights[dst_offset] = (sum / 4.0);
+    let probabilities = (vec4<f32>(a, (a + b), ((a + b) + c), sum) / max(sum, 0.0001));
+    textureStore(tex_out, vec2<i32>(coord.xy), probabilities);
+  }
+}
diff --git a/test/benchmark/simple_compute.wgsl b/test/benchmark/simple_compute.wgsl
new file mode 100644
index 0000000..f5b883a
--- /dev/null
+++ b/test/benchmark/simple_compute.wgsl
@@ -0,0 +1,10 @@
+struct SB {
+  data : array<i32>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> buffer : SB;
+
+[[stage(compute), workgroup_size(1, 2, 3)]]
+fn main([[builtin(global_invocation_id)]] id : vec3<u32>) {
+  buffer.data[id.x] = buffer.data[id.x] + 1;
+}
diff --git a/test/benchmark/simple_compute.wgsl.expected.hlsl b/test/benchmark/simple_compute.wgsl.expected.hlsl
new file mode 100644
index 0000000..92d72b0
--- /dev/null
+++ b/test/benchmark/simple_compute.wgsl.expected.hlsl
@@ -0,0 +1,15 @@
+RWByteAddressBuffer buffer : register(u0, space0);
+
+struct tint_symbol_1 {
+  uint3 id : SV_DispatchThreadID;
+};
+
+void main_inner(uint3 id) {
+  buffer.Store((4u * id.x), asuint((asint(buffer.Load((4u * id.x))) + 1)));
+}
+
+[numthreads(1, 2, 3)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.id);
+  return;
+}
diff --git a/test/benchmark/simple_compute.wgsl.expected.msl b/test/benchmark/simple_compute.wgsl.expected.msl
new file mode 100644
index 0000000..a2119e6
--- /dev/null
+++ b/test/benchmark/simple_compute.wgsl.expected.msl
@@ -0,0 +1,16 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct SB {
+  /* 0x0000 */ int data[1];
+};
+
+void tint_symbol_1_inner(uint3 id, device SB* const tint_symbol_2) {
+  (*(tint_symbol_2)).data[id[0]] = as_type<int>((as_type<uint>((*(tint_symbol_2)).data[id[0]]) + as_type<uint>(1)));
+}
+
+kernel void tint_symbol_1(device SB* tint_symbol_3 [[buffer(0)]], uint3 id [[thread_position_in_grid]]) {
+  tint_symbol_1_inner(id, tint_symbol_3);
+  return;
+}
+
diff --git a/test/benchmark/simple_compute.wgsl.expected.spvasm b/test/benchmark/simple_compute.wgsl.expected.spvasm
new file mode 100644
index 0000000..35c2109
--- /dev/null
+++ b/test/benchmark/simple_compute.wgsl.expected.spvasm
@@ -0,0 +1,55 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %id_1
+               OpExecutionMode %main LocalSize 1 2 3
+               OpName %id_1 "id_1"
+               OpName %SB "SB"
+               OpMemberName %SB 0 "data"
+               OpName %buffer "buffer"
+               OpName %main_inner "main_inner"
+               OpName %id "id"
+               OpName %main "main"
+               OpDecorate %id_1 BuiltIn GlobalInvocationId
+               OpDecorate %SB Block
+               OpMemberDecorate %SB 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %buffer DescriptorSet 0
+               OpDecorate %buffer Binding 0
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+       %id_1 = OpVariable %_ptr_Input_v3uint Input
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+         %SB = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_SB = OpTypePointer StorageBuffer %SB
+     %buffer = OpVariable %_ptr_StorageBuffer_SB StorageBuffer
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void %v3uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+      %int_1 = OpConstant %int 1
+         %24 = OpTypeFunction %void
+ %main_inner = OpFunction %void None %10
+         %id = OpFunctionParameter %v3uint
+         %14 = OpLabel
+         %16 = OpCompositeExtract %uint %id 0
+         %18 = OpAccessChain %_ptr_StorageBuffer_int %buffer %uint_0 %16
+         %19 = OpCompositeExtract %uint %id 0
+         %20 = OpAccessChain %_ptr_StorageBuffer_int %buffer %uint_0 %19
+         %21 = OpLoad %int %20
+         %23 = OpIAdd %int %21 %int_1
+               OpStore %18 %23
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %24
+         %26 = OpLabel
+         %28 = OpLoad %v3uint %id_1
+         %27 = OpFunctionCall %void %main_inner %28
+               OpReturn
+               OpFunctionEnd
diff --git a/test/benchmark/simple_compute.wgsl.expected.wgsl b/test/benchmark/simple_compute.wgsl.expected.wgsl
new file mode 100644
index 0000000..ac98410
--- /dev/null
+++ b/test/benchmark/simple_compute.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+struct SB {
+  data : array<i32>;
+};
+
+[[group(0), binding(0)]] var<storage, read_write> buffer : SB;
+
+[[stage(compute), workgroup_size(1, 2, 3)]]
+fn main([[builtin(global_invocation_id)]] id : vec3<u32>) {
+  buffer.data[id.x] = (buffer.data[id.x] + 1);
+}
diff --git a/test/benchmark/simple_fragment.wgsl b/test/benchmark/simple_fragment.wgsl
new file mode 100644
index 0000000..bb5ed94
--- /dev/null
+++ b/test/benchmark/simple_fragment.wgsl
@@ -0,0 +1,12 @@
+struct Input {
+  [[location(0)]] color: vec4<f32>;
+};
+
+struct Output {
+  [[location(0)]] color: vec4<f32>;
+};
+
+[[stage(fragment)]]
+fn main(in : Input) -> Output {
+  return Output(in.color);
+}
diff --git a/test/benchmark/simple_fragment.wgsl.expected.hlsl b/test/benchmark/simple_fragment.wgsl.expected.hlsl
new file mode 100644
index 0000000..4a462c8
--- /dev/null
+++ b/test/benchmark/simple_fragment.wgsl.expected.hlsl
@@ -0,0 +1,25 @@
+struct Input {
+  float4 color;
+};
+struct Output {
+  float4 color;
+};
+struct tint_symbol_2 {
+  float4 color : TEXCOORD0;
+};
+struct tint_symbol_3 {
+  float4 color : SV_Target0;
+};
+
+Output main_inner(Input tint_symbol) {
+  const Output tint_symbol_4 = {tint_symbol.color};
+  return tint_symbol_4;
+}
+
+tint_symbol_3 main(tint_symbol_2 tint_symbol_1) {
+  const Input tint_symbol_5 = {tint_symbol_1.color};
+  const Output inner_result = main_inner(tint_symbol_5);
+  tint_symbol_3 wrapper_result = (tint_symbol_3)0;
+  wrapper_result.color = inner_result.color;
+  return wrapper_result;
+}
diff --git a/test/benchmark/simple_fragment.wgsl.expected.msl b/test/benchmark/simple_fragment.wgsl.expected.msl
new file mode 100644
index 0000000..0bc1116
--- /dev/null
+++ b/test/benchmark/simple_fragment.wgsl.expected.msl
@@ -0,0 +1,29 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Input {
+  float4 color;
+};
+struct Output {
+  float4 color;
+};
+struct tint_symbol_2 {
+  float4 color [[user(locn0)]];
+};
+struct tint_symbol_3 {
+  float4 color [[color(0)]];
+};
+
+Output tint_symbol_inner(Input in) {
+  Output const tint_symbol_4 = {.color=in.color};
+  return tint_symbol_4;
+}
+
+fragment tint_symbol_3 tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  Input const tint_symbol_5 = {.color=tint_symbol_1.color};
+  Output const inner_result = tint_symbol_inner(tint_symbol_5);
+  tint_symbol_3 wrapper_result = {};
+  wrapper_result.color = inner_result.color;
+  return wrapper_result;
+}
+
diff --git a/test/benchmark/simple_fragment.wgsl.expected.spvasm b/test/benchmark/simple_fragment.wgsl.expected.spvasm
new file mode 100644
index 0000000..531ed06
--- /dev/null
+++ b/test/benchmark/simple_fragment.wgsl.expected.spvasm
@@ -0,0 +1,50 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main" %color_1 %color_2
+               OpExecutionMode %main OriginUpperLeft
+               OpName %color_1 "color_1"
+               OpName %color_2 "color_2"
+               OpName %Output "Output"
+               OpMemberName %Output 0 "color"
+               OpName %Input "Input"
+               OpMemberName %Input 0 "color"
+               OpName %main_inner "main_inner"
+               OpName %in "in"
+               OpName %main "main"
+               OpDecorate %color_1 Location 0
+               OpDecorate %color_2 Location 0
+               OpMemberDecorate %Output 0 Offset 0
+               OpMemberDecorate %Input 0 Offset 0
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+    %color_1 = OpVariable %_ptr_Input_v4float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %7 = OpConstantNull %v4float
+    %color_2 = OpVariable %_ptr_Output_v4float Output %7
+     %Output = OpTypeStruct %v4float
+      %Input = OpTypeStruct %v4float
+          %8 = OpTypeFunction %Output %Input
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+ %main_inner = OpFunction %Output None %8
+         %in = OpFunctionParameter %Input
+         %13 = OpLabel
+         %14 = OpCompositeExtract %v4float %in 0
+         %15 = OpCompositeConstruct %Output %14
+               OpReturnValue %15
+               OpFunctionEnd
+       %main = OpFunction %void None %16
+         %19 = OpLabel
+         %21 = OpLoad %v4float %color_1
+         %22 = OpCompositeConstruct %Input %21
+         %20 = OpFunctionCall %Output %main_inner %22
+         %23 = OpCompositeExtract %v4float %20 0
+               OpStore %color_2 %23
+               OpReturn
+               OpFunctionEnd
diff --git a/test/benchmark/simple_fragment.wgsl.expected.wgsl b/test/benchmark/simple_fragment.wgsl.expected.wgsl
new file mode 100644
index 0000000..1f4225e
--- /dev/null
+++ b/test/benchmark/simple_fragment.wgsl.expected.wgsl
@@ -0,0 +1,14 @@
+struct Input {
+  [[location(0)]]
+  color : vec4<f32>;
+};
+
+struct Output {
+  [[location(0)]]
+  color : vec4<f32>;
+};
+
+[[stage(fragment)]]
+fn main(in : Input) -> Output {
+  return Output(in.color);
+}
diff --git a/test/benchmark/simple_vertex.wgsl b/test/benchmark/simple_vertex.wgsl
new file mode 100644
index 0000000..8f9a07b
--- /dev/null
+++ b/test/benchmark/simple_vertex.wgsl
@@ -0,0 +1,12 @@
+struct Input {
+  [[location(0)]] position: vec4<f32>;
+};
+
+struct Output {
+  [[builtin(position)]] position : vec4<f32>;
+};
+
+[[stage(vertex)]]
+fn main(in : Input) -> Output {
+  return Output(in.position);
+}
diff --git a/test/benchmark/simple_vertex.wgsl.expected.hlsl b/test/benchmark/simple_vertex.wgsl.expected.hlsl
new file mode 100644
index 0000000..03b0d54
--- /dev/null
+++ b/test/benchmark/simple_vertex.wgsl.expected.hlsl
@@ -0,0 +1,25 @@
+struct Input {
+  float4 position;
+};
+struct Output {
+  float4 position;
+};
+struct tint_symbol_2 {
+  float4 position : TEXCOORD0;
+};
+struct tint_symbol_3 {
+  float4 position : SV_Position;
+};
+
+Output main_inner(Input tint_symbol) {
+  const Output tint_symbol_4 = {tint_symbol.position};
+  return tint_symbol_4;
+}
+
+tint_symbol_3 main(tint_symbol_2 tint_symbol_1) {
+  const Input tint_symbol_5 = {tint_symbol_1.position};
+  const Output inner_result = main_inner(tint_symbol_5);
+  tint_symbol_3 wrapper_result = (tint_symbol_3)0;
+  wrapper_result.position = inner_result.position;
+  return wrapper_result;
+}
diff --git a/test/benchmark/simple_vertex.wgsl.expected.msl b/test/benchmark/simple_vertex.wgsl.expected.msl
new file mode 100644
index 0000000..3c18b5f
--- /dev/null
+++ b/test/benchmark/simple_vertex.wgsl.expected.msl
@@ -0,0 +1,29 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Input {
+  float4 position;
+};
+struct Output {
+  float4 position;
+};
+struct tint_symbol_2 {
+  float4 position [[attribute(0)]];
+};
+struct tint_symbol_3 {
+  float4 position [[position]];
+};
+
+Output tint_symbol_inner(Input in) {
+  Output const tint_symbol_4 = {.position=in.position};
+  return tint_symbol_4;
+}
+
+vertex tint_symbol_3 tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  Input const tint_symbol_5 = {.position=tint_symbol_1.position};
+  Output const inner_result = tint_symbol_inner(tint_symbol_5);
+  tint_symbol_3 wrapper_result = {};
+  wrapper_result.position = inner_result.position;
+  return wrapper_result;
+}
+
diff --git a/test/benchmark/simple_vertex.wgsl.expected.spvasm b/test/benchmark/simple_vertex.wgsl.expected.spvasm
new file mode 100644
index 0000000..6490ac0
--- /dev/null
+++ b/test/benchmark/simple_vertex.wgsl.expected.spvasm
@@ -0,0 +1,56 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 28
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Vertex %main "main" %position_1 %position_2 %vertex_point_size
+               OpName %position_1 "position_1"
+               OpName %position_2 "position_2"
+               OpName %vertex_point_size "vertex_point_size"
+               OpName %Output "Output"
+               OpMemberName %Output 0 "position"
+               OpName %Input "Input"
+               OpMemberName %Input 0 "position"
+               OpName %main_inner "main_inner"
+               OpName %in "in"
+               OpName %main "main"
+               OpDecorate %position_1 Location 0
+               OpDecorate %position_2 BuiltIn Position
+               OpDecorate %vertex_point_size BuiltIn PointSize
+               OpMemberDecorate %Output 0 Offset 0
+               OpMemberDecorate %Input 0 Offset 0
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%_ptr_Input_v4float = OpTypePointer Input %v4float
+ %position_1 = OpVariable %_ptr_Input_v4float Input
+%_ptr_Output_v4float = OpTypePointer Output %v4float
+          %7 = OpConstantNull %v4float
+ %position_2 = OpVariable %_ptr_Output_v4float Output %7
+%_ptr_Output_float = OpTypePointer Output %float
+         %10 = OpConstantNull %float
+%vertex_point_size = OpVariable %_ptr_Output_float Output %10
+     %Output = OpTypeStruct %v4float
+      %Input = OpTypeStruct %v4float
+         %11 = OpTypeFunction %Output %Input
+       %void = OpTypeVoid
+         %19 = OpTypeFunction %void
+    %float_1 = OpConstant %float 1
+ %main_inner = OpFunction %Output None %11
+         %in = OpFunctionParameter %Input
+         %16 = OpLabel
+         %17 = OpCompositeExtract %v4float %in 0
+         %18 = OpCompositeConstruct %Output %17
+               OpReturnValue %18
+               OpFunctionEnd
+       %main = OpFunction %void None %19
+         %22 = OpLabel
+         %24 = OpLoad %v4float %position_1
+         %25 = OpCompositeConstruct %Input %24
+         %23 = OpFunctionCall %Output %main_inner %25
+         %26 = OpCompositeExtract %v4float %23 0
+               OpStore %position_2 %26
+               OpStore %vertex_point_size %float_1
+               OpReturn
+               OpFunctionEnd
diff --git a/test/benchmark/simple_vertex.wgsl.expected.wgsl b/test/benchmark/simple_vertex.wgsl.expected.wgsl
new file mode 100644
index 0000000..7bc18b1
--- /dev/null
+++ b/test/benchmark/simple_vertex.wgsl.expected.wgsl
@@ -0,0 +1,14 @@
+struct Input {
+  [[location(0)]]
+  position : vec4<f32>;
+};
+
+struct Output {
+  [[builtin(position)]]
+  position : vec4<f32>;
+};
+
+[[stage(vertex)]]
+fn main(in : Input) -> Output {
+  return Output(in.position);
+}
diff --git a/third_party/CMakeLists.txt b/third_party/CMakeLists.txt
index f54ed0a..47d6664 100644
--- a/third_party/CMakeLists.txt
+++ b/third_party/CMakeLists.txt
@@ -12,6 +12,11 @@
 # See the License for the specific language governing permissions and
 # limitations under the License.
 
+if (${TINT_BUILD_BENCHMARKS})
+  set(BENCHMARK_ENABLE_TESTING FALSE CACHE BOOL FALSE FORCE)
+  add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/benchmark EXCLUDE_FROM_ALL)
+endif()
+
 if (${TINT_BUILD_TESTS} AND NOT TARGET gmock)
   set(gtest_force_shared_crt ON CACHE BOOL "Controls whether a shared run-time library should be used even when Google Test is built as static library" FORCE)
   add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/googletest EXCLUDE_FROM_ALL)
diff --git a/tools/benchdiff b/tools/benchdiff
new file mode 100755
index 0000000..3cc175f
--- /dev/null
+++ b/tools/benchdiff
@@ -0,0 +1,33 @@
+#!/usr/bin/env bash
+# Copyright 2022 The Tint Authors
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+set -e # Fail on any error.
+
+if [ ! -x "$(which go)" ] ; then
+    echo "error: go needs to be on \$PATH to use $0"
+    exit 1
+fi
+
+SCRIPT_DIR="$( cd "$( dirname "${BASH_SOURCE[0]}")" >/dev/null 2>&1 && pwd )"
+ROOT_DIR="$( cd "${SCRIPT_DIR}/.." >/dev/null 2>&1 && pwd )"
+BINARY="${SCRIPT_DIR}/bin/benchdiff"
+
+# Rebuild the binary.
+# Note, go caches build artifacts, so this is quick for repeat calls
+pushd "${SCRIPT_DIR}/src/cmd/benchdiff" > /dev/null
+    go build -o "${BINARY}" main.go
+popd > /dev/null
+
+"${BINARY}" "$@"
diff --git a/tools/src/bench/bench.go b/tools/src/bench/bench.go
new file mode 100644
index 0000000..8bb1d99
--- /dev/null
+++ b/tools/src/bench/bench.go
@@ -0,0 +1,150 @@
+// Copyright 2022 The Tint Authors
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+// Package bench provides types and methods for parsing Google benchmark results.
+package bench
+
+import (
+	"encoding/json"
+	"errors"
+	"fmt"
+	"regexp"
+	"strconv"
+	"strings"
+	"time"
+)
+
+// Test holds the results of a single benchmark test.
+type Test struct {
+	Name       string
+	NumTasks   uint
+	NumThreads uint
+	Duration   time.Duration
+	Iterations uint
+}
+
+var testVarRE = regexp.MustCompile(`([\w])+:([0-9]+)`)
+
+func (t *Test) parseName() {
+	for _, match := range testVarRE.FindAllStringSubmatch(t.Name, -1) {
+		if len(match) != 3 {
+			continue
+		}
+		n, err := strconv.Atoi(match[2])
+		if err != nil {
+			continue
+		}
+		switch match[1] {
+		case "threads":
+			t.NumThreads = uint(n)
+		case "tasks":
+			t.NumTasks = uint(n)
+		}
+	}
+}
+
+// Benchmark holds a set of benchmark test results.
+type Benchmark struct {
+	Tests []Test
+}
+
+// Parse parses the benchmark results from the string s.
+// Parse will handle the json and 'console' formats.
+func Parse(s string) (Benchmark, error) {
+	type Parser = func(s string) (Benchmark, error)
+	for _, parser := range []Parser{parseConsole, parseJSON} {
+		b, err := parser(s)
+		switch err {
+		case nil:
+			return b, nil
+		case errWrongFormat:
+		default:
+			return Benchmark{}, err
+		}
+	}
+
+	return Benchmark{}, errors.New("Unrecognised file format")
+}
+
+var errWrongFormat = errors.New("Wrong format")
+var consoleLineRE = regexp.MustCompile(`([\w/:]+)\s+([0-9]+(?:.[0-9]+)?) ns\s+[0-9]+(?:.[0-9]+) ns\s+([0-9]+)`)
+
+func parseConsole(s string) (Benchmark, error) {
+	blocks := strings.Split(s, "------------------------------------------------------------------------------------------")
+	if len(blocks) != 3 {
+		return Benchmark{}, errWrongFormat
+	}
+
+	lines := strings.Split(blocks[2], "\n")
+	b := Benchmark{
+		Tests: make([]Test, 0, len(lines)),
+	}
+	for _, line := range lines {
+		if len(line) == 0 {
+			continue
+		}
+		matches := consoleLineRE.FindStringSubmatch(line)
+		if len(matches) != 4 {
+			return Benchmark{}, fmt.Errorf("Unable to parse the line:\n" + line)
+		}
+		ns, err := strconv.ParseFloat(matches[2], 64)
+		if err != nil {
+			return Benchmark{}, fmt.Errorf("Unable to parse the duration: " + matches[2])
+		}
+		iterations, err := strconv.Atoi(matches[3])
+		if err != nil {
+			return Benchmark{}, fmt.Errorf("Unable to parse the number of iterations: " + matches[3])
+		}
+
+		t := Test{
+			Name:       matches[1],
+			Duration:   time.Nanosecond * time.Duration(ns),
+			Iterations: uint(iterations),
+		}
+		t.parseName()
+		b.Tests = append(b.Tests, t)
+	}
+	return b, nil
+}
+
+func parseJSON(s string) (Benchmark, error) {
+	type T struct {
+		Name       string  `json:"name"`
+		Iterations uint    `json:"iterations"`
+		Time       float64 `json:"real_time"`
+	}
+	type B struct {
+		Tests []T `json:"benchmarks"`
+	}
+	b := B{}
+	d := json.NewDecoder(strings.NewReader(s))
+	if err := d.Decode(&b); err != nil {
+		return Benchmark{}, err
+	}
+
+	out := Benchmark{
+		Tests: make([]Test, len(b.Tests)),
+	}
+	for i, test := range b.Tests {
+		t := Test{
+			Name:       test.Name,
+			Duration:   time.Nanosecond * time.Duration(int64(test.Time)),
+			Iterations: test.Iterations,
+		}
+		t.parseName()
+		out.Tests[i] = t
+	}
+
+	return out, nil
+}
diff --git a/tools/src/cmd/benchdiff/main.go b/tools/src/cmd/benchdiff/main.go
new file mode 100644
index 0000000..14ec4f3
--- /dev/null
+++ b/tools/src/cmd/benchdiff/main.go
@@ -0,0 +1,177 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+// benchdiff is a tool that compares two Google benchmark results and displays
+// sorted performance differences.
+package main
+
+import (
+	"errors"
+	"flag"
+	"fmt"
+	"io/ioutil"
+	"os"
+	"path/filepath"
+	"sort"
+	"strings"
+	"text/tabwriter"
+	"time"
+
+	"dawn.googlesource.com/tint/tools/src/bench"
+)
+
+var (
+	minDiff    = flag.Duration("min-diff", time.Microsecond*10, "Filter away time diffs less than this duration")
+	minRelDiff = flag.Float64("min-rel-diff", 0.01, "Filter away absolute relative diffs between [1, 1+x]")
+)
+
+func main() {
+	flag.ErrHelp = errors.New("benchdiff is a tool to compare two benchmark results")
+	flag.Parse()
+	flag.Usage = func() {
+		fmt.Fprintln(os.Stderr, "benchdiff <benchmark-a> <benchmark-b>")
+		flag.PrintDefaults()
+	}
+
+	args := flag.Args()
+	if len(args) < 2 {
+		flag.Usage()
+		os.Exit(1)
+	}
+
+	pathA, pathB := args[0], args[1]
+
+	if err := run(pathA, pathB); err != nil {
+		fmt.Fprintln(os.Stderr, err)
+		os.Exit(-1)
+	}
+}
+
+func run(pathA, pathB string) error {
+	fileA, err := ioutil.ReadFile(pathA)
+	if err != nil {
+		return err
+	}
+	benchA, err := bench.Parse(string(fileA))
+	if err != nil {
+		return err
+	}
+
+	fileB, err := ioutil.ReadFile(pathB)
+	if err != nil {
+		return err
+	}
+	benchB, err := bench.Parse(string(fileB))
+	if err != nil {
+		return err
+	}
+
+	compare(benchA, benchB, fileName(pathA), fileName(pathB))
+
+	return nil
+}
+
+func fileName(path string) string {
+	_, name := filepath.Split(path)
+	return name
+}
+
+func compare(benchA, benchB bench.Benchmark, nameA, nameB string) {
+	type times struct {
+		a time.Duration
+		b time.Duration
+	}
+	byName := map[string]times{}
+	for _, test := range benchA.Tests {
+		byName[test.Name] = times{a: test.Duration}
+	}
+	for _, test := range benchB.Tests {
+		t := byName[test.Name]
+		t.b = test.Duration
+		byName[test.Name] = t
+	}
+
+	type delta struct {
+		name       string
+		times      times
+		relDiff    float64
+		absRelDiff float64
+	}
+	deltas := []delta{}
+	for name, times := range byName {
+		if times.a == 0 || times.b == 0 {
+			continue // Assuming test was missing from a or b
+		}
+		diff := times.b - times.a
+		absDiff := diff
+		if absDiff < 0 {
+			absDiff = -absDiff
+		}
+		if absDiff < *minDiff {
+			continue
+		}
+
+		relDiff := float64(times.b) / float64(times.a)
+		absRelDiff := relDiff
+		if absRelDiff < 1 {
+			absRelDiff = 1.0 / absRelDiff
+		}
+		if absRelDiff < (1.0 + *minRelDiff) {
+			continue
+		}
+
+		d := delta{
+			name:       name,
+			times:      times,
+			relDiff:    relDiff,
+			absRelDiff: absRelDiff,
+		}
+		deltas = append(deltas, d)
+	}
+
+	sort.Slice(deltas, func(i, j int) bool { return deltas[j].relDiff < deltas[i].relDiff })
+
+	fmt.Println("A:", nameA)
+	fmt.Println("B:", nameB)
+	fmt.Println()
+
+	buf := strings.Builder{}
+	{
+		w := tabwriter.NewWriter(&buf, 1, 1, 0, ' ', 0)
+		fmt.Fprintln(w, "Test name\t | Δ (A → B)\t | % (A → B)\t | % (B → A)\t | × (A → B)\t | × (B → A)\t | A \t | B")
+		fmt.Fprintln(w, "\t-+\t-+\t-+\t-+\t-+\t-+\t-+\t-")
+		for _, delta := range deltas {
+			a2b := delta.times.b - delta.times.a
+			fmt.Fprintf(w, "%v \t | %v \t | %+2.1f%% \t | %+2.1f%% \t | %+.4f \t | %+.4f \t | %v \t | %v \t|\n",
+				delta.name,
+				a2b, // Δ (A → B)
+				100*float64(a2b)/float64(delta.times.a),       // % (A → B)
+				100*float64(-a2b)/float64(delta.times.b),      // % (B → A)
+				float64(delta.times.b)/float64(delta.times.a), // × (A → B)
+				float64(delta.times.a)/float64(delta.times.b), // × (B → A)
+				delta.times.a, // A
+				delta.times.b, // B
+			)
+		}
+		w.Flush()
+	}
+
+	// Split the table by line so we can add in a header line
+	lines := strings.Split(buf.String(), "\n")
+	fmt.Println(lines[0])
+	fmt.Println(strings.ReplaceAll(lines[1], " ", "-"))
+	for _, l := range lines[2:] {
+		fmt.Println(l)
+	}
+}