diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 7e85a38..47efd52 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -359,10 +359,7 @@
     "transform/transform.cc",
     "transform/transform.h",
   ]
-  deps = [
-    ":libtint_program_src",
-    ":libtint_utils_src",
-  ]
+  deps = [ ":libtint_utils_src" ]
 }
 
 libtint_source_set("libtint_ast_transform_base_src") {
@@ -370,11 +367,11 @@
     "ast/transform/transform.cc",
     "ast/transform/transform.h",
   ]
-  public_deps = [ ":libtint_transform_src" ]
   deps = [
     ":libtint_builtins_src",
     ":libtint_program_src",
     ":libtint_sem_src",
+    ":libtint_transform_src",
     ":libtint_type_src",
     ":libtint_utils_src",
   ]
@@ -387,7 +384,10 @@
   ]
   deps = [
     ":libtint_ast_transform_base_src",
+    ":libtint_ir_builder_src",
+    ":libtint_ir_src",
     ":libtint_program_src",
+    ":libtint_transform_src",
   ]
 }
 
@@ -492,9 +492,9 @@
     "ast/transform/zero_init_workgroup_memory.cc",
     "ast/transform/zero_init_workgroup_memory.h",
   ]
-  public_deps = [ ":libtint_ast_transform_base_src" ]
   deps = [
     ":libtint_ast_src",
+    ":libtint_ast_transform_base_src",
     ":libtint_builtins_src",
     ":libtint_program_src",
     ":libtint_sem_src",
@@ -505,6 +505,22 @@
   ]
 }
 
+if (tint_build_ir) {
+  libtint_source_set("libtint_ir_transform_src") {
+    sources = [
+      "ir/transform/add_empty_entry_point.cc",
+      "ir/transform/add_empty_entry_point.h",
+    ]
+    deps = [
+      ":libtint_builtins_src",
+      ":libtint_ir_src",
+      ":libtint_symbols_src",
+      ":libtint_type_src",
+      ":libtint_utils_src",
+    ]
+  }
+}
+
 libtint_source_set("libtint_ast_hdrs") {
   sources = [
     "ast/accessor_expression.h",
@@ -1014,12 +1030,13 @@
 
   if (tint_build_ir) {
     sources += [
-      "writer/spirv/generator_impl_ir.cc",
-      "writer/spirv/generator_impl_ir.h",
+      "writer/spirv/ir/generator_impl_ir.cc",
+      "writer/spirv/ir/generator_impl_ir.h",
     ]
     deps += [
       ":libtint_ir_builder_src",
       ":libtint_ir_src",
+      ":libtint_ir_transform_src",
     ]
   }
 }
@@ -1185,6 +1202,8 @@
     "ir/bitcast.h",
     "ir/block.cc",
     "ir/block.h",
+    "ir/block_param.cc",
+    "ir/block_param.h",
     "ir/builder.cc",
     "ir/builder.h",
     "ir/builtin.cc",
@@ -1213,6 +1232,8 @@
     "ir/if.h",
     "ir/instruction.cc",
     "ir/instruction.h",
+    "ir/load.cc",
+    "ir/load.h",
     "ir/loop.cc",
     "ir/loop.h",
     "ir/module.cc",
@@ -1223,6 +1244,8 @@
     "ir/store.h",
     "ir/switch.cc",
     "ir/switch.h",
+    "ir/transform/transform.cc",
+    "ir/transform/transform.h",
     "ir/unary.cc",
     "ir/unary.h",
     "ir/user_call.cc",
@@ -1237,6 +1260,7 @@
     ":libtint_builtins_src",
     ":libtint_constant_src",
     ":libtint_symbols_src",
+    ":libtint_transform_src",
     ":libtint_type_src",
     ":libtint_utils_src",
   ]
@@ -1245,6 +1269,7 @@
 source_set("libtint") {
   public_deps = [
     ":libtint_ast_src",
+    ":libtint_ast_transform_base_src",
     ":libtint_ast_transform_src",
     ":libtint_constant_src",
     ":libtint_initializer_src",
@@ -1736,6 +1761,7 @@
     ]
 
     deps = [
+      ":libtint_ast_transform_base_src",
       ":libtint_ast_transform_src",
       ":libtint_builtins_src",
       ":libtint_transform_manager_src",
@@ -1746,6 +1772,21 @@
     ]
   }
 
+  if (tint_build_ir) {
+    tint_unittests_source_set("tint_unittests_ir_transform_src") {
+      sources = [
+        "ir/transform/add_empty_entry_point_test.cc",
+        "ir/transform/test_helper.h",
+      ]
+
+      deps = [
+        ":libtint_ir_src",
+        ":libtint_ir_transform_src",
+        ":libtint_transform_manager_src",
+      ]
+    }
+  }
+
   tint_unittests_source_set("tint_unittests_utils_src") {
     sources = [
       "debug_test.cc",
@@ -1902,11 +1943,14 @@
 
     if (tint_build_ir) {
       sources += [
-        "writer/spirv/generator_impl_binary_test.cc",
-        "writer/spirv/generator_impl_constant_test.cc",
-        "writer/spirv/generator_impl_function_test.cc",
-        "writer/spirv/generator_impl_ir_test.cc",
-        "writer/spirv/generator_impl_type_test.cc",
+        "writer/spirv/ir/generator_impl_ir_binary_test.cc",
+        "writer/spirv/ir/generator_impl_ir_constant_test.cc",
+        "writer/spirv/ir/generator_impl_ir_function_test.cc",
+        "writer/spirv/ir/generator_impl_ir_if_test.cc",
+        "writer/spirv/ir/generator_impl_ir_test.cc",
+        "writer/spirv/ir/generator_impl_ir_type_test.cc",
+        "writer/spirv/ir/generator_impl_ir_var_test.cc",
+        "writer/spirv/ir/test_helper_ir.h",
       ]
       deps += [ ":libtint_ir_src" ]
     }
@@ -2181,12 +2225,23 @@
       "clone_context_test.cc",
       "program_builder_test.cc",
       "program_test.cc",
+      "transform/manager_test.cc",
     ]
 
     deps = [
+      ":libtint_ast_transform_base_src",
+      ":libtint_program_src",
+      ":libtint_transform_manager_src",
       ":libtint_unittests_ast_helper",
       ":tint_unittests_ast_src",
     ]
+
+    if (tint_build_ir) {
+      deps += [
+        ":libtint_ir_builder_src",
+        ":libtint_ir_src",
+      ]
+    }
   }
 
   tint_unittests_source_set("tint_unittests_ir_src") {
@@ -2203,6 +2258,7 @@
       "ir/from_program_test.cc",
       "ir/from_program_unary_test.cc",
       "ir/from_program_var_test.cc",
+      "ir/load_test.cc",
       "ir/module_test.cc",
       "ir/store_test.cc",
       "ir/test_helper.h",
@@ -2278,7 +2334,10 @@
     }
 
     if (tint_build_ir) {
-      deps += [ ":tint_unittests_ir_src" ]
+      deps += [
+        ":tint_unittests_ir_src",
+        ":tint_unittests_ir_transform_src",
+      ]
     }
 
     if (build_with_chromium) {
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 7d12ae2..d0c7487 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -656,8 +656,8 @@
 
   if(${TINT_BUILD_IR})
     list(APPEND TINT_LIB_SRCS
-      writer/spirv/generator_impl_ir.cc
-      writer/spirv/generator_impl_ir.h
+      writer/spirv/ir/generator_impl_ir.cc
+      writer/spirv/ir/generator_impl_ir.h
     )
   endif()
 endif()
@@ -716,6 +716,8 @@
     ir/bitcast.h
     ir/block.cc
     ir/block.h
+    ir/block_param.cc
+    ir/block_param.h
     ir/builder.cc
     ir/builder.h
     ir/builtin.cc
@@ -746,6 +748,8 @@
     ir/if.h
     ir/instruction.cc
     ir/instruction.h
+    ir/load.cc
+    ir/load.h
     ir/loop.cc
     ir/loop.h
     ir/module.cc
@@ -766,6 +770,10 @@
     ir/value.h
     ir/var.cc
     ir/var.h
+    ir/transform/add_empty_entry_point.cc
+    ir/transform/add_empty_entry_point.h
+    ir/transform/transform.cc
+    ir/transform/transform.h
   )
 endif()
 
@@ -996,6 +1004,7 @@
     symbol_test.cc
     test_main.cc
     ast/transform/transform_test.cc
+    transform/manager_test.cc
     type/array_test.cc
     type/atomic_test.cc
     type/bool_test.cc
@@ -1231,12 +1240,14 @@
 
     if(${TINT_BUILD_IR})
       list(APPEND TINT_TEST_SRCS
-        writer/spirv/generator_impl_binary_test.cc
-        writer/spirv/generator_impl_constant_test.cc
-        writer/spirv/generator_impl_function_test.cc
-        writer/spirv/generator_impl_ir_test.cc
-        writer/spirv/generator_impl_type_test.cc
-        writer/spirv/test_helper_ir.h
+        writer/spirv/ir/generator_impl_ir_binary_test.cc
+        writer/spirv/ir/generator_impl_ir_constant_test.cc
+        writer/spirv/ir/generator_impl_ir_function_test.cc
+        writer/spirv/ir/generator_impl_ir_if_test.cc
+        writer/spirv/ir/generator_impl_ir_test.cc
+        writer/spirv/ir/generator_impl_ir_type_test.cc
+        writer/spirv/ir/generator_impl_ir_var_test.cc
+        writer/spirv/ir/test_helper_ir.h
       )
     endif()
   endif()
@@ -1459,6 +1470,7 @@
       ir/from_program_test.cc
       ir/from_program_unary_test.cc
       ir/from_program_var_test.cc
+      ir/load_test.cc
       ir/module_test.cc
       ir/store_test.cc
       ir/test_helper.h
diff --git a/src/tint/ir/block.h b/src/tint/ir/block.h
index 3981355..5af32fe 100644
--- a/src/tint/ir/block.h
+++ b/src/tint/ir/block.h
@@ -15,6 +15,7 @@
 #ifndef SRC_TINT_IR_BLOCK_H_
 #define SRC_TINT_IR_BLOCK_H_
 
+#include "src/tint/ir/block_param.h"
 #include "src/tint/ir/branch.h"
 #include "src/tint/ir/flow_node.h"
 #include "src/tint/ir/instruction.h"
@@ -40,6 +41,9 @@
 
     /// The instructions in the block
     utils::Vector<const Instruction*, 16> instructions;
+
+    /// The parameters passed into the block
+    utils::Vector<const BlockParam*, 0> params;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/block_param.cc b/src/tint/ir/block_param.cc
new file mode 100644
index 0000000..f014d19
--- /dev/null
+++ b/src/tint/ir/block_param.cc
@@ -0,0 +1,25 @@
+// Copyright 2023 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/tint/ir/block_param.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::BlockParam);
+
+namespace tint::ir {
+
+BlockParam::BlockParam(const type::Type* ty) : type(ty) {}
+
+BlockParam::~BlockParam() = default;
+
+}  // namespace tint::ir
diff --git a/src/tint/ir/block_param.h b/src/tint/ir/block_param.h
new file mode 100644
index 0000000..8ba68a7
--- /dev/null
+++ b/src/tint/ir/block_param.h
@@ -0,0 +1,45 @@
+// Copyright 2023 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_TINT_IR_BLOCK_PARAM_H_
+#define SRC_TINT_IR_BLOCK_PARAM_H_
+
+#include "src/tint/ir/value.h"
+#include "src/tint/utils/castable.h"
+
+namespace tint::ir {
+
+/// An instruction in the IR.
+class BlockParam : public utils::Castable<BlockParam, Value> {
+  public:
+    /// Constructor
+    /// @param type the type of the var
+    explicit BlockParam(const type::Type* type);
+    BlockParam(const BlockParam& inst) = delete;
+    BlockParam(BlockParam&& inst) = delete;
+    ~BlockParam() override;
+
+    BlockParam& operator=(const BlockParam& inst) = delete;
+    BlockParam& operator=(BlockParam&& inst) = delete;
+
+    /// @returns the type of the var
+    const type::Type* Type() const override { return type; }
+
+    /// the result type of the instruction
+    const type::Type* type = nullptr;
+};
+
+}  // namespace tint::ir
+
+#endif  // SRC_TINT_IR_BLOCK_PARAM_H_
diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index d62bb43..983815e 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -17,6 +17,8 @@
 #include <utility>
 
 #include "src/tint/constant/scalar.h"
+#include "src/tint/type/pointer.h"
+#include "src/tint/type/reference.h"
 
 namespace tint::ir {
 
@@ -114,7 +116,7 @@
 }
 
 Binary* Builder::CreateBinary(Binary::Kind kind, const type::Type* type, Value* lhs, Value* rhs) {
-    return ir.instructions.Create<ir::Binary>(kind, type, lhs, rhs);
+    return ir.values.Create<ir::Binary>(kind, type, lhs, rhs);
 }
 
 Binary* Builder::And(const type::Type* type, Value* lhs, Value* rhs) {
@@ -182,21 +184,13 @@
 }
 
 Unary* Builder::CreateUnary(Unary::Kind kind, const type::Type* type, Value* val) {
-    return ir.instructions.Create<ir::Unary>(kind, type, val);
-}
-
-Unary* Builder::AddressOf(const type::Type* type, Value* val) {
-    return CreateUnary(Unary::Kind::kAddressOf, type, val);
+    return ir.values.Create<ir::Unary>(kind, type, val);
 }
 
 Unary* Builder::Complement(const type::Type* type, Value* val) {
     return CreateUnary(Unary::Kind::kComplement, type, val);
 }
 
-Unary* Builder::Indirection(const type::Type* type, Value* val) {
-    return CreateUnary(Unary::Kind::kIndirection, type, val);
-}
-
 Unary* Builder::Negation(const type::Type* type, Value* val) {
     return CreateUnary(Unary::Kind::kNegation, type, val);
 }
@@ -206,43 +200,51 @@
 }
 
 ir::Bitcast* Builder::Bitcast(const type::Type* type, Value* val) {
-    return ir.instructions.Create<ir::Bitcast>(type, val);
+    return ir.values.Create<ir::Bitcast>(type, val);
 }
 
 ir::Discard* Builder::Discard() {
-    return ir.instructions.Create<ir::Discard>();
+    return ir.values.Create<ir::Discard>();
 }
 
 ir::UserCall* Builder::UserCall(const type::Type* type,
                                 Symbol name,
                                 utils::VectorRef<Value*> args) {
-    return ir.instructions.Create<ir::UserCall>(type, name, std::move(args));
+    return ir.values.Create<ir::UserCall>(type, name, std::move(args));
 }
 
 ir::Convert* Builder::Convert(const type::Type* to,
                               const type::Type* from,
                               utils::VectorRef<Value*> args) {
-    return ir.instructions.Create<ir::Convert>(to, from, std::move(args));
+    return ir.values.Create<ir::Convert>(to, from, std::move(args));
 }
 
 ir::Construct* Builder::Construct(const type::Type* to, utils::VectorRef<Value*> args) {
-    return ir.instructions.Create<ir::Construct>(to, std::move(args));
+    return ir.values.Create<ir::Construct>(to, std::move(args));
 }
 
 ir::Builtin* Builder::Builtin(const type::Type* type,
                               builtin::Function func,
                               utils::VectorRef<Value*> args) {
-    return ir.instructions.Create<ir::Builtin>(type, func, args);
+    return ir.values.Create<ir::Builtin>(type, func, args);
+}
+
+ir::Load* Builder::Load(Value* from) {
+    auto* ptr = from->Type()->As<type::Pointer>();
+    TINT_ASSERT(IR, ptr);
+    return ir.values.Create<ir::Load>(ptr->StoreType(), from);
 }
 
 ir::Store* Builder::Store(Value* to, Value* from) {
-    return ir.instructions.Create<ir::Store>(to, from);
+    return ir.values.Create<ir::Store>(to, from);
 }
 
-ir::Var* Builder::Declare(const type::Type* type,
-                          builtin::AddressSpace address_space,
-                          builtin::Access access) {
-    return ir.instructions.Create<ir::Var>(type, address_space, access);
+ir::Var* Builder::Declare(const type::Type* type) {
+    return ir.values.Create<ir::Var>(type);
+}
+
+ir::BlockParam* Builder::BlockParam(const type::Type* type) {
+    return ir.values.Create<ir::BlockParam>(type);
 }
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index b4a2a81..7c0cab6 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -28,6 +28,7 @@
 #include "src/tint/ir/function.h"
 #include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
+#include "src/tint/ir/load.h"
 #include "src/tint/ir/loop.h"
 #include "src/tint/ir/module.h"
 #include "src/tint/ir/root_terminator.h"
@@ -279,24 +280,12 @@
     /// @returns the operation
     Unary* CreateUnary(Unary::Kind kind, const type::Type* type, Value* val);
 
-    /// Creates an AddressOf operation
-    /// @param type the result type of the expression
-    /// @param val the value
-    /// @returns the operation
-    Unary* AddressOf(const type::Type* type, Value* val);
-
     /// Creates a Complement operation
     /// @param type the result type of the expression
     /// @param val the value
     /// @returns the operation
     Unary* Complement(const type::Type* type, Value* val);
 
-    /// Creates an Indirection operation
-    /// @param type the result type of the expression
-    /// @param val the value
-    /// @returns the operation
-    Unary* Indirection(const type::Type* type, Value* val);
-
     /// Creates a Negation operation
     /// @param type the result type of the expression
     /// @param val the value
@@ -350,7 +339,12 @@
                          builtin::Function func,
                          utils::VectorRef<Value*> args);
 
-    /// Creates an store instruction
+    /// Creates a load instruction
+    /// @param from the expression being loaded from
+    /// @returns the instruction
+    ir::Load* Load(Value* from);
+
+    /// Creates a store instruction
     /// @param to the expression being stored too
     /// @param from the expression being stored
     /// @returns the instruction
@@ -358,12 +352,13 @@
 
     /// Creates a new `var` declaration
     /// @param type the var type
-    /// @param address_space the address space
-    /// @param access the access mode
     /// @returns the instruction
-    ir::Var* Declare(const type::Type* type,
-                     builtin::AddressSpace address_space,
-                     builtin::Access access);
+    ir::Var* Declare(const type::Type* type);
+
+    /// Creates a new `BlockParam`
+    /// @param type the parameter type
+    /// @returns the value
+    ir::BlockParam* BlockParam(const type::Type* type);
 
     /// Retrieves the root block for the module, creating if necessary
     /// @returns the root block
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index d4ee3a7..f68b21a 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -27,6 +27,7 @@
 #include "src/tint/ir/discard.h"
 #include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
+#include "src/tint/ir/load.h"
 #include "src/tint/ir/loop.h"
 #include "src/tint/ir/root_terminator.h"
 #include "src/tint/ir/store.h"
@@ -154,7 +155,19 @@
                 return;
             }
 
-            Indent() << "%fn" << IdOf(b) << " = block {" << std::endl;
+            Indent() << "%fn" << IdOf(b) << " = block";
+            if (!b->params.IsEmpty()) {
+                out_ << " (";
+                for (auto* p : b->params) {
+                    if (p != b->params.Front()) {
+                        out_ << ", ";
+                    }
+                    EmitValue(p);
+                }
+                out_ << ")";
+            }
+
+            out_ << " {" << std::endl;
             {
                 ScopedIndent si(indent_size_);
                 EmitBlockInstructions(b);
@@ -248,7 +261,20 @@
         [&](const ir::If* i) {
             Indent() << "%fn" << IdOf(i) << " = if ";
             EmitValue(i->condition);
-            out_ << " [t: %fn" << IdOf(i->true_.target) << ", f: %fn" << IdOf(i->false_.target);
+
+            bool has_true = !i->true_.target->IsDead();
+            bool has_false = !i->false_.target->IsDead();
+
+            out_ << " [";
+            if (has_true) {
+                out_ << "t: %fn" << IdOf(i->true_.target);
+            }
+            if (has_false) {
+                if (has_true) {
+                    out_ << ", ";
+                }
+                out_ << "f: %fn" << IdOf(i->false_.target);
+            }
             if (i->merge.target->IsConnected()) {
                 out_ << ", m: %fn" << IdOf(i->merge.target);
             }
@@ -258,10 +284,12 @@
                 ScopedIndent if_indent(indent_size_);
                 ScopedStopNode scope(stop_nodes_, i->merge.target);
 
-                Indent() << "# true branch" << std::endl;
-                Walk(i->true_.target);
+                if (has_true) {
+                    Indent() << "# true branch" << std::endl;
+                    Walk(i->true_.target);
+                }
 
-                if (!i->false_.target->IsDead()) {
+                if (has_false) {
                     Indent() << "# false branch" << std::endl;
                     Walk(i->false_.target);
                 }
@@ -324,6 +352,13 @@
     return out_.str();
 }
 
+void Disassembler::EmitValueWithType(const Value* val) {
+    EmitValue(val);
+    if (auto* i = val->As<ir::Instruction>(); i->Type() != nullptr) {
+        out_ << ":" << i->Type()->FriendlyName();
+    }
+}
+
 void Disassembler::EmitValue(const Value* val) {
     tint::Switch(
         val,
@@ -368,12 +403,11 @@
             };
             emit(constant->value);
         },
-        [&](const ir::Instruction* i) {
-            out_ << "%" << IdOf(i);
-            if (i->Type() != nullptr) {
-                out_ << ":" << i->Type()->FriendlyName();
-            }
-        });
+        [&](const ir::Instruction* i) { out_ << "%" << IdOf(i); },
+        [&](const ir::BlockParam* p) {
+            out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
+        },
+        [&](Default) { out_ << "Unknown value: " << val->TypeInfo().name; });
 }
 
 void Disassembler::EmitInstruction(const Instruction* inst) {
@@ -381,26 +415,31 @@
         inst,  //
         [&](const ir::Binary* b) { EmitBinary(b); }, [&](const ir::Unary* u) { EmitUnary(u); },
         [&](const ir::Bitcast* b) {
-            EmitValue(b);
+            EmitValueWithType(b);
             out_ << " = bitcast ";
             EmitArgs(b);
         },
         [&](const ir::Discard*) { out_ << "discard"; },
         [&](const ir::Builtin* b) {
-            EmitValue(b);
+            EmitValueWithType(b);
             out_ << " = " << builtin::str(b->Func()) << " ";
             EmitArgs(b);
         },
         [&](const ir::Construct* c) {
-            EmitValue(c);
+            EmitValueWithType(c);
             out_ << " = construct ";
             EmitArgs(c);
         },
         [&](const ir::Convert* c) {
-            EmitValue(c);
+            EmitValueWithType(c);
             out_ << " = convert " << c->FromType()->FriendlyName() << ", ";
             EmitArgs(c);
         },
+        [&](const ir::Load* l) {
+            EmitValueWithType(l);
+            out_ << " = load ";
+            EmitValue(l->from);
+        },
         [&](const ir::Store* s) {
             out_ << "store ";
             EmitValue(s->to);
@@ -408,7 +447,7 @@
             EmitValue(s->from);
         },
         [&](const ir::UserCall* uc) {
-            EmitValue(uc);
+            EmitValueWithType(uc);
             out_ << " = call " << uc->name.Name();
             if (uc->args.Length() > 0) {
                 out_ << ", ";
@@ -416,8 +455,8 @@
             EmitArgs(uc);
         },
         [&](const ir::Var* v) {
-            EmitValue(v);
-            out_ << " = var " << v->address_space << ", " << v->access;
+            EmitValueWithType(v);
+            out_ << " = var";
             if (v->initializer) {
                 out_ << ", ";
                 EmitValue(v->initializer);
@@ -437,7 +476,7 @@
 }
 
 void Disassembler::EmitBinary(const Binary* b) {
-    EmitValue(b);
+    EmitValueWithType(b);
     out_ << " = ";
     switch (b->kind) {
         case Binary::Kind::kAdd:
@@ -496,18 +535,12 @@
 }
 
 void Disassembler::EmitUnary(const Unary* u) {
-    EmitValue(u);
+    EmitValueWithType(u);
     out_ << " = ";
     switch (u->kind) {
-        case Unary::Kind::kAddressOf:
-            out_ << "addr_of";
-            break;
         case Unary::Kind::kComplement:
             out_ << "complement";
             break;
-        case Unary::Kind::kIndirection:
-            out_ << "indirection";
-            break;
         case Unary::Kind::kNegation:
             out_ << "negation";
             break;
diff --git a/src/tint/ir/disassembler.h b/src/tint/ir/disassembler.h
index 2438d80..c8953db 100644
--- a/src/tint/ir/disassembler.h
+++ b/src/tint/ir/disassembler.h
@@ -55,6 +55,7 @@
 
     void Walk(const FlowNode* node);
     void EmitInstruction(const Instruction* inst);
+    void EmitValueWithType(const Value* val);
     void EmitValue(const Value* val);
     void EmitArgs(const Call* call);
     void EmitBinary(const Binary* b);
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index f7729f3..bb9c3f9 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -48,6 +48,7 @@
 #include "src/tint/ast/literal_expression.h"
 #include "src/tint/ast/loop_statement.h"
 #include "src/tint/ast/override.h"
+#include "src/tint/ast/phony_expression.h"
 #include "src/tint/ast/return_statement.h"
 #include "src/tint/ast/statement.h"
 #include "src/tint/ast/struct.h"
@@ -59,6 +60,7 @@
 #include "src/tint/ast/var.h"
 #include "src/tint/ast/variable_decl_statement.h"
 #include "src/tint/ast/while_statement.h"
+#include "src/tint/ir/block_param.h"
 #include "src/tint/ir/builder.h"
 #include "src/tint/ir/function.h"
 #include "src/tint/ir/if.h"
@@ -72,6 +74,7 @@
 #include "src/tint/sem/builtin.h"
 #include "src/tint/sem/call.h"
 #include "src/tint/sem/function.h"
+#include "src/tint/sem/load.h"
 #include "src/tint/sem/materialize.h"
 #include "src/tint/sem/module.h"
 #include "src/tint/sem/switch_statement.h"
@@ -80,6 +83,8 @@
 #include "src/tint/sem/value_expression.h"
 #include "src/tint/sem/variable.h"
 #include "src/tint/switch.h"
+#include "src/tint/type/pointer.h"
+#include "src/tint/type/reference.h"
 #include "src/tint/type/void.h"
 #include "src/tint/utils/defer.h"
 #include "src/tint/utils/result.h"
@@ -336,7 +341,6 @@
             current_flow_block_ = ir_func->start_target;
             EmitBlock(ast_func->body);
 
-            // TODO(dsinclair): Store return type and attributes
             // TODO(dsinclair): Store parameters
 
             // If the branch target has already been set then a `return` was called. Only set in the
@@ -390,6 +394,16 @@
     }
 
     void EmitAssignment(const ast::AssignmentStatement* stmt) {
+        // If assigning to a phony, just generate the RHS and we're done. Note that, because this
+        // isn't used, a subsequent transform could remove it due to it being dead code. This could
+        // then change the interface for the program (i.e. a global var no longer used). If that
+        // happens we have to either fix this to store to a phony value, or make sure we pull the
+        // interface before doing the dead code elimination.
+        if (stmt->lhs->Is<ast::PhonyExpression>()) {
+            (void)EmitExpression(stmt->rhs);
+            return;
+        }
+
         auto lhs = EmitExpression(stmt->lhs);
         if (!lhs) {
             return;
@@ -409,15 +423,20 @@
             return;
         }
 
-        auto* ty = lhs.Get()->Type();
-        auto* rhs = ty->UnwrapRef()->is_signed_integer_scalar() ? builder_.Constant(1_i)
-                                                                : builder_.Constant(1_u);
+        // Load from the LHS.
+        auto* lhs_value = builder_.Load(lhs.Get());
+        current_flow_block_->instructions.Push(lhs_value);
+
+        auto* ty = lhs_value->Type();
+
+        auto* rhs =
+            ty->is_signed_integer_scalar() ? builder_.Constant(1_i) : builder_.Constant(1_u);
 
         Binary* inst = nullptr;
         if (stmt->increment) {
-            inst = builder_.Add(ty, lhs.Get(), rhs);
+            inst = builder_.Add(ty, lhs_value, rhs);
         } else {
-            inst = builder_.Subtract(ty, lhs.Get(), rhs);
+            inst = builder_.Subtract(ty, lhs_value, rhs);
         }
         current_flow_block_->instructions.Push(inst);
 
@@ -435,38 +454,44 @@
         if (!rhs) {
             return;
         }
-        auto* ty = lhs.Get()->Type();
+
+        // Load from the LHS.
+        auto* lhs_value = builder_.Load(lhs.Get());
+        current_flow_block_->instructions.Push(lhs_value);
+
+        auto* ty = lhs_value->Type();
+
         Binary* inst = nullptr;
         switch (stmt->op) {
             case ast::BinaryOp::kAnd:
-                inst = builder_.And(ty, lhs.Get(), rhs.Get());
+                inst = builder_.And(ty, lhs_value, rhs.Get());
                 break;
             case ast::BinaryOp::kOr:
-                inst = builder_.Or(ty, lhs.Get(), rhs.Get());
+                inst = builder_.Or(ty, lhs_value, rhs.Get());
                 break;
             case ast::BinaryOp::kXor:
-                inst = builder_.Xor(ty, lhs.Get(), rhs.Get());
+                inst = builder_.Xor(ty, lhs_value, rhs.Get());
                 break;
             case ast::BinaryOp::kShiftLeft:
-                inst = builder_.ShiftLeft(ty, lhs.Get(), rhs.Get());
+                inst = builder_.ShiftLeft(ty, lhs_value, rhs.Get());
                 break;
             case ast::BinaryOp::kShiftRight:
-                inst = builder_.ShiftRight(ty, lhs.Get(), rhs.Get());
+                inst = builder_.ShiftRight(ty, lhs_value, rhs.Get());
                 break;
             case ast::BinaryOp::kAdd:
-                inst = builder_.Add(ty, lhs.Get(), rhs.Get());
+                inst = builder_.Add(ty, lhs_value, rhs.Get());
                 break;
             case ast::BinaryOp::kSubtract:
-                inst = builder_.Subtract(ty, lhs.Get(), rhs.Get());
+                inst = builder_.Subtract(ty, lhs_value, rhs.Get());
                 break;
             case ast::BinaryOp::kMultiply:
-                inst = builder_.Multiply(ty, lhs.Get(), rhs.Get());
+                inst = builder_.Multiply(ty, lhs_value, rhs.Get());
                 break;
             case ast::BinaryOp::kDivide:
-                inst = builder_.Divide(ty, lhs.Get(), rhs.Get());
+                inst = builder_.Divide(ty, lhs_value, rhs.Get());
                 break;
             case ast::BinaryOp::kModulo:
-                inst = builder_.Modulo(ty, lhs.Get(), rhs.Get());
+                inst = builder_.Modulo(ty, lhs_value, rhs.Get());
                 break;
             case ast::BinaryOp::kLessThanEqual:
             case ast::BinaryOp::kGreaterThanEqual:
@@ -796,7 +821,8 @@
 
     utils::Result<Value*> EmitExpression(const ast::Expression* expr) {
         // If this is a value that has been const-eval'd return the result.
-        if (auto* sem = program_->Sem().Get(expr)->As<sem::ValueExpression>()) {
+        auto* sem = program_->Sem().GetVal(expr);
+        if (sem) {
             if (auto* v = sem->ConstantValue()) {
                 if (auto* cv = v->Clone(clone_ctx_)) {
                     return builder_.Constant(cv);
@@ -804,7 +830,7 @@
             }
         }
 
-        return tint::Switch(
+        auto result = tint::Switch(
             expr,
             // [&](const ast::IndexAccessorExpression* a) {
             // TODO(dsinclair): Implement
@@ -825,15 +851,23 @@
             // [&](const ast::MemberAccessorExpression* m) {
             // TODO(dsinclair): Implement
             // },
-            // [&](const ast::PhonyExpression*) {
-            // TODO(dsinclair): Implement. The call may have side effects so has to be made.
-            // },
             [&](const ast::UnaryOpExpression* u) { return EmitUnary(u); },
+            // Note, ast::PhonyExpression is explicitly not handled here as it should never get into
+            // this method. The assignment statement should have filtered it out already.
             [&](Default) {
                 add_error(expr->source,
                           "unknown expression type: " + std::string(expr->TypeInfo().name));
                 return utils::Failure;
             });
+
+        // If this expression maps to sem::Load, insert a load instruction to get the result.
+        if (result && sem->Is<sem::Load>()) {
+            auto* load = builder_.Load(result.Get());
+            current_flow_block_->instructions.Push(load);
+            return load;
+        }
+
+        return result;
     }
 
     void EmitVariable(const ast::Variable* var) {
@@ -842,8 +876,12 @@
         return tint::Switch(  //
             var,
             [&](const ast::Var* v) {
-                auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx);
-                auto* val = builder_.Declare(ty, sem->AddressSpace(), sem->Access());
+                auto* ref = sem->Type()->As<type::Reference>();
+                auto* ty = builder_.ir.types.Get<type::Pointer>(
+                    ref->StoreType()->Clone(clone_ctx_.type_ctx), ref->AddressSpace(),
+                    ref->Access());
+
+                auto* val = builder_.Declare(ty);
                 current_flow_block_->instructions.Push(val);
 
                 if (v->initializer) {
@@ -904,14 +942,12 @@
         Instruction* inst = nullptr;
         switch (expr->op) {
             case ast::UnaryOp::kAddressOf:
-                inst = builder_.AddressOf(ty, val.Get());
-                break;
+            case ast::UnaryOp::kIndirection:
+                // 'address-of' and 'indirection' just fold away and we propagate the pointer.
+                return val;
             case ast::UnaryOp::kComplement:
                 inst = builder_.Complement(ty, val.Get());
                 break;
-            case ast::UnaryOp::kIndirection:
-                inst = builder_.Indirection(ty, val.Get());
-                break;
             case ast::UnaryOp::kNegation:
                 inst = builder_.Negation(ty, val.Get());
                 break;
@@ -943,27 +979,34 @@
             return utils::Failure;
         }
 
-        // Generate a variable to store the short-circut into
-        auto* ty = builder_.ir.types.Get<type::Bool>();
-        auto* result_var =
-            builder_.Declare(ty, builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
-        current_flow_block_->instructions.Push(result_var);
-
-        auto* lhs_store = builder_.Store(result_var, lhs.Get());
-        current_flow_block_->instructions.Push(lhs_store);
-
         auto* if_node = builder_.CreateIf(lhs.Get());
         BranchTo(if_node);
 
+        auto* result = builder_.BlockParam(builder_.ir.types.Get<type::Bool>());
+        if_node->merge.target->As<Block>()->params.Push(result);
+
         utils::Result<Value*> rhs;
         {
             FlowStackScope scope(this, if_node);
 
+            utils::Vector<Value*, 1> alt_args;
+            alt_args.Push(lhs.Get());
+
             // If this is an `&&` then we only evaluate the RHS expression in the true block.
             // If this is an `||` then we only evaluate the RHS expression in the false block.
             if (expr->op == ast::BinaryOp::kLogicalAnd) {
+                // If the lhs is false, then that is the result we want to pass to the merge block
+                // as our argument
+                current_flow_block_ = if_node->false_.target->As<Block>();
+                BranchTo(if_node->merge.target, std::move(alt_args));
+
                 current_flow_block_ = if_node->true_.target->As<Block>();
             } else {
+                // If the lhs is true, then that is the result we want to pass to the merge block
+                // as our argument
+                current_flow_block_ = if_node->true_.target->As<Block>();
+                BranchTo(if_node->merge.target, std::move(alt_args));
+
                 current_flow_block_ = if_node->false_.target->As<Block>();
             }
 
@@ -971,14 +1014,14 @@
             if (!rhs) {
                 return utils::Failure;
             }
-            auto* rhs_store = builder_.Store(result_var, rhs.Get());
-            current_flow_block_->instructions.Push(rhs_store);
+            utils::Vector<Value*, 1> args;
+            args.Push(rhs.Get());
 
-            BranchTo(if_node->merge.target);
+            BranchTo(if_node->merge.target, std::move(args));
         }
         current_flow_block_ = if_node->merge.target->As<Block>();
 
-        return result_var;
+        return result;
     }
 
     utils::Result<Value*> EmitBinary(const ast::BinaryExpression* expr) {
diff --git a/src/tint/ir/from_program_binary_test.cc b/src/tint/ir/from_program_binary_test.cc
index 33266cb..85377ff 100644
--- a/src/tint/ir/from_program_binary_test.cc
+++ b/src/tint/ir/from_program_binary_test.cc
@@ -42,7 +42,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:u32 = add %1:u32, 4u
+    %tint_symbol:u32 = add %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -58,14 +58,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var private, read_write
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = add %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:u32 = load %v1
+    %3:u32 = add %2, 1u
+    store %v1, %3
   } -> %func_end # return
 } %func_end
 
@@ -81,14 +82,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var private, read_write
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = add %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:u32 = load %v1
+    %3:u32 = add %2, 1u
+    store %v1, %3
   } -> %func_end # return
 } %func_end
 
@@ -111,7 +113,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:u32 = sub %1:u32, 4u
+    %tint_symbol:u32 = sub %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -127,14 +129,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, i32, read_write> = var private, read_write
+  %v1:ptr<private, i32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, i32, read_write> = sub %v1:ref<private, i32, read_write>, 1i
-    store %v1:ref<private, i32, read_write>, %2:ref<private, i32, read_write>
+    %2:i32 = load %v1
+    %3:i32 = sub %2, 1i
+    store %v1, %3
   } -> %func_end # return
 } %func_end
 
@@ -150,14 +153,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var private, read_write
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = sub %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:u32 = load %v1
+    %3:u32 = sub %2, 1u
+    store %v1, %3
   } -> %func_end # return
 } %func_end
 
@@ -180,7 +184,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:u32 = mul %1:u32, 4u
+    %tint_symbol:u32 = mul %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -196,14 +200,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var private, read_write
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = mul %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:u32 = load %v1
+    %3:u32 = mul %2, 1u
+    store %v1, %3
   } -> %func_end # return
 } %func_end
 
@@ -226,7 +231,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:u32 = div %1:u32, 4u
+    %tint_symbol:u32 = div %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -242,14 +247,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var private, read_write
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = div %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:u32 = load %v1
+    %3:u32 = div %2, 1u
+    store %v1, %3
   } -> %func_end # return
 } %func_end
 
@@ -272,7 +278,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:u32 = mod %1:u32, 4u
+    %tint_symbol:u32 = mod %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -288,14 +294,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var private, read_write
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = mod %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:u32 = load %v1
+    %3:u32 = mod %2, 1u
+    store %v1, %3
   } -> %func_end # return
 } %func_end
 
@@ -318,7 +325,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:u32 = and %1:u32, 4u
+    %tint_symbol:u32 = and %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -334,14 +341,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, bool, read_write> = var private, read_write
+  %v1:ptr<private, bool, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, bool, read_write> = and %v1:ref<private, bool, read_write>, false
-    store %v1:ref<private, bool, read_write>, %2:ref<private, bool, read_write>
+    %2:bool = load %v1
+    %3:bool = and %2, false
+    store %v1, %3
   } -> %func_end # return
 } %func_end
 
@@ -364,7 +372,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:u32 = or %1:u32, 4u
+    %tint_symbol:u32 = or %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -380,14 +388,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, bool, read_write> = var private, read_write
+  %v1:ptr<private, bool, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, bool, read_write> = or %v1:ref<private, bool, read_write>, false
-    store %v1:ref<private, bool, read_write>, %2:ref<private, bool, read_write>
+    %2:bool = load %v1
+    %3:bool = or %2, false
+    store %v1, %3
   } -> %func_end # return
 } %func_end
 
@@ -410,7 +419,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:u32 = xor %1:u32, 4u
+    %tint_symbol:u32 = xor %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -426,14 +435,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var private, read_write
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = xor %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:u32 = load %v1
+    %3:u32 = xor %2, 1u
+    store %v1, %3
   } -> %func_end # return
 } %func_end
 
@@ -442,7 +452,7 @@
 
 TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
     Func("my_func", utils::Empty, ty.bool_(), utils::Vector{Return(true)});
-    auto* expr = LogicalAnd(Call("my_func"), false);
+    auto* expr = If(LogicalAnd(Call("my_func"), false), Block());
     WrapInFunction(expr);
 
     auto m = Build();
@@ -456,18 +466,32 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:bool = call my_func
-    %tint_symbol:bool = var function, read_write
-    store %tint_symbol:bool, %1:bool
   } -> %fn5 # branch
 
-  %fn5 = if %1:bool [t: %fn6, f: %fn7, m: %fn8]
+  %fn5 = if %1 [t: %fn6, f: %fn7, m: %fn8]
     # true branch
     %fn6 = block {
-      store %tint_symbol:bool, false
-    } -> %fn8 # branch
+    } -> %fn8 false # branch
+
+    # false branch
+    %fn7 = block {
+    } -> %fn8 %1 # branch
 
   # if merge
-  %fn8 = block {
+  %fn8 = block (%2:bool) {
+  } -> %fn9 # branch
+
+  %fn9 = if %2:bool [t: %fn10, f: %fn11, m: %fn12]
+    # true branch
+    %fn10 = block {
+    } -> %fn12 # branch
+
+    # false branch
+    %fn11 = block {
+    } -> %fn12 # branch
+
+  # if merge
+  %fn12 = block {
   } -> %func_end # return
 } %func_end
 
@@ -476,7 +500,7 @@
 
 TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
     Func("my_func", utils::Empty, ty.bool_(), utils::Vector{Return(true)});
-    auto* expr = LogicalOr(Call("my_func"), true);
+    auto* expr = If(LogicalOr(Call("my_func"), true), Block());
     WrapInFunction(expr);
 
     auto m = Build();
@@ -490,19 +514,32 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:bool = call my_func
-    %tint_symbol:bool = var function, read_write
-    store %tint_symbol:bool, %1:bool
   } -> %fn5 # branch
 
-  %fn5 = if %1:bool [t: %fn6, f: %fn7, m: %fn8]
+  %fn5 = if %1 [t: %fn6, f: %fn7, m: %fn8]
     # true branch
+    %fn6 = block {
+    } -> %fn8 %1 # branch
+
     # false branch
     %fn7 = block {
-      store %tint_symbol:bool, true
-    } -> %fn8 # branch
+    } -> %fn8 true # branch
 
   # if merge
-  %fn8 = block {
+  %fn8 = block (%2:bool) {
+  } -> %fn9 # branch
+
+  %fn9 = if %2:bool [t: %fn10, f: %fn11, m: %fn12]
+    # true branch
+    %fn10 = block {
+    } -> %fn12 # branch
+
+    # false branch
+    %fn11 = block {
+    } -> %fn12 # branch
+
+  # if merge
+  %fn12 = block {
   } -> %func_end # return
 } %func_end
 
@@ -525,7 +562,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:bool = eq %1:u32, 4u
+    %tint_symbol:bool = eq %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -548,7 +585,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:bool = neq %1:u32, 4u
+    %tint_symbol:bool = neq %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -571,7 +608,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:bool = lt %1:u32, 4u
+    %tint_symbol:bool = lt %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -594,7 +631,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:bool = gt %1:u32, 4u
+    %tint_symbol:bool = gt %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -617,7 +654,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:bool = lte %1:u32, 4u
+    %tint_symbol:bool = lte %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -640,7 +677,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:bool = gte %1:u32, 4u
+    %tint_symbol:bool = gte %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -663,7 +700,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:u32 = shiftl %1:u32, 4u
+    %tint_symbol:u32 = shiftl %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -679,14 +716,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var private, read_write
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = shiftl %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:u32 = load %v1
+    %3:u32 = shiftl %2, 1u
+    store %v1, %3
   } -> %func_end # return
 } %func_end
 
@@ -709,7 +747,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:u32 = shiftr %1:u32, 4u
+    %tint_symbol:u32 = shiftr %1, 4u
   } -> %func_end # return
 } %func_end
 
@@ -725,14 +763,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var private, read_write
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = shiftr %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:u32 = load %v1
+    %3:u32 = shiftr %2, 1u
+    store %v1, %3
   } -> %func_end # return
 } %func_end
 
@@ -757,24 +796,25 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:f32 = call my_func
-    %2:bool = lt %1:f32, 2.0f
-    %tint_symbol:bool = var function, read_write
-    store %tint_symbol:bool, %2:bool
+    %2:bool = lt %1, 2.0f
   } -> %fn5 # branch
 
-  %fn5 = if %2:bool [t: %fn6, f: %fn7, m: %fn8]
+  %fn5 = if %2 [t: %fn6, f: %fn7, m: %fn8]
     # true branch
     %fn6 = block {
+      %3:f32 = call my_func
       %4:f32 = call my_func
-      %5:f32 = call my_func
-      %6:f32 = mul 2.29999995231628417969f, %5:f32
-      %7:f32 = div %4:f32, %6:f32
-      %8:bool = gt 2.5f, %7:f32
-      store %tint_symbol:bool, %8:bool
-    } -> %fn8 # branch
+      %5:f32 = mul 2.29999995231628417969f, %4
+      %6:f32 = div %3, %5
+      %7:bool = gt 2.5f, %6
+    } -> %fn8 %7 # branch
+
+    # false branch
+    %fn7 = block {
+    } -> %fn8 %2 # branch
 
   # if merge
-  %fn8 = block {
+  %fn8 = block (%tint_symbol:bool) {
   } -> %func_end # return
 } %func_end
 
diff --git a/src/tint/ir/from_program_call_test.cc b/src/tint/ir/from_program_call_test.cc
index ae217d1..038f4dd 100644
--- a/src/tint/ir/from_program_call_test.cc
+++ b/src/tint/ir/from_program_call_test.cc
@@ -43,7 +43,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:f32 = call my_func
-    %tint_symbol:f32 = bitcast %1:f32
+    %tint_symbol:f32 = bitcast %1
   } -> %func_end # return
 } %func_end
 
@@ -100,13 +100,14 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %i:ref<private, i32, read_write> = var private, read_write, 1i
+  %i:ptr<private, i32, read_write> = var, 1i
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %tint_symbol:f32 = convert i32, %i:ref<private, i32, read_write>
+    %2:i32 = load %i
+    %tint_symbol:f32 = convert i32, %2
   } -> %func_end # return
 } %func_end
 
@@ -121,7 +122,7 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %i:ref<private, vec3<f32>, read_write> = var private, read_write, vec3<f32> 0.0f
+  %i:ptr<private, vec3<f32>, read_write> = var, vec3<f32> 0.0f
 }
 
 
@@ -137,13 +138,14 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %i:ref<private, f32, read_write> = var private, read_write, 1.0f
+  %i:ptr<private, f32, read_write> = var, 1.0f
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %i:ref<private, f32, read_write>
+    %2:f32 = load %i
+    %tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %2
   } -> %func_end # return
 } %func_end
 
diff --git a/src/tint/ir/from_program_store_test.cc b/src/tint/ir/from_program_store_test.cc
index 4e89dbe..5bb6398 100644
--- a/src/tint/ir/from_program_store_test.cc
+++ b/src/tint/ir/from_program_store_test.cc
@@ -36,13 +36,13 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %a:ref<private, u32, read_write> = var private, read_write
+  %a:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    store %a:ref<private, u32, read_write>, 4u
+    store %a, 4u
   } -> %func_end # return
 } %func_end
 
diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc
index 234b6cc..47dc911 100644
--- a/src/tint/ir/from_program_test.cc
+++ b/src/tint/ir/from_program_test.cc
@@ -1427,5 +1427,27 @@
 )");
 }
 
+TEST_F(IR_BuilderImplTest, Emit_Phony) {
+    Func("b", utils::Empty, ty.i32(), Return(1_i));
+    WrapInFunction(Ignore(Call("b")));
+
+    auto m = Build();
+    ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+    EXPECT_EQ(Disassemble(m.Get()),
+              R"(%fn1 = func b():i32 {
+  %fn2 = block {
+  } -> %func_end 1i # return
+} %func_end
+
+%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+  %fn4 = block {
+    %1:i32 = call b
+  } -> %func_end # return
+} %func_end
+
+)");
+}
+
 }  // namespace
 }  // namespace tint::ir
diff --git a/src/tint/ir/from_program_unary_test.cc b/src/tint/ir/from_program_unary_test.cc
index b2ba2d7..2afd2a2 100644
--- a/src/tint/ir/from_program_unary_test.cc
+++ b/src/tint/ir/from_program_unary_test.cc
@@ -42,7 +42,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:bool = call my_func
-    %tint_symbol:bool = eq %1:bool, false
+    %tint_symbol:bool = eq %1, false
   } -> %func_end # return
 } %func_end
 
@@ -65,7 +65,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:u32 = call my_func
-    %tint_symbol:u32 = complement %1:u32
+    %tint_symbol:u32 = complement %1
   } -> %func_end # return
 } %func_end
 
@@ -88,7 +88,7 @@
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
     %1:i32 = call my_func
-    %tint_symbol:i32 = negation %1:i32
+    %tint_symbol:i32 = negation %1
   } -> %func_end # return
 } %func_end
 
@@ -105,13 +105,12 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, i32, read_write> = var private, read_write
+  %v2:ptr<private, i32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %v2:ptr<private, i32, read_write> = addr_of %v1:ref<private, i32, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -122,7 +121,7 @@
     GlobalVar("v1", builtin::AddressSpace::kPrivate, ty.i32());
     utils::Vector stmts = {
         Decl(Let("v3", AddressOf("v1"))),
-        Decl(Let("v2", Deref("v3"))),
+        Assign(Deref("v3"), 42_i),
     };
     WrapInFunction(stmts);
 
@@ -130,14 +129,13 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, i32, read_write> = var private, read_write
+  %v3:ptr<private, i32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %v3:ptr<private, i32, read_write> = addr_of %v1:ref<private, i32, read_write>
-    %v2:i32 = indirection %v3:ptr<private, i32, read_write>
+    store %v3, 42i
   } -> %func_end # return
 } %func_end
 
diff --git a/src/tint/ir/from_program_var_test.cc b/src/tint/ir/from_program_var_test.cc
index b5379ac..e235f88 100644
--- a/src/tint/ir/from_program_var_test.cc
+++ b/src/tint/ir/from_program_var_test.cc
@@ -33,7 +33,7 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %a:ref<private, u32, read_write> = var private, read_write
+  %a:ptr<private, u32, read_write> = var
 }
 
 
@@ -48,7 +48,7 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %a:ref<private, u32, read_write> = var private, read_write, 2u
+  %a:ptr<private, u32, read_write> = var, 2u
 }
 
 
@@ -65,7 +65,7 @@
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn2 = block {
-    %a:ref<function, u32, read_write> = var function, read_write
+    %a:ptr<function, u32, read_write> = var
   } -> %func_end # return
 } %func_end
 
@@ -83,7 +83,7 @@
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn2 = block {
-    %a:ref<function, u32, read_write> = var function, read_write, 2u
+    %a:ptr<function, u32, read_write> = var, 2u
   } -> %func_end # return
 } %func_end
 
diff --git a/src/tint/writer/spirv/generator_impl_ir_test.cc b/src/tint/ir/load.cc
similarity index 60%
copy from src/tint/writer/spirv/generator_impl_ir_test.cc
copy to src/tint/ir/load.cc
index a202eea..1fe55c0 100644
--- a/src/tint/writer/spirv/generator_impl_ir_test.cc
+++ b/src/tint/ir/load.cc
@@ -12,18 +12,19 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "src/tint/writer/spirv/test_helper_ir.h"
+#include "src/tint/ir/load.h"
+#include "src/tint/debug.h"
 
-namespace tint::writer::spirv {
-namespace {
+TINT_INSTANTIATE_TYPEINFO(tint::ir::Load);
 
-TEST_F(SpvGeneratorImplTest, ModuleHeader) {
-    ASSERT_TRUE(generator_.Generate()) << generator_.Diagnostics().str();
-    auto got = Disassemble(generator_.Result());
-    EXPECT_EQ(got, R"(OpCapability Shader
-OpMemoryModel Logical GLSL450
-)");
+namespace tint::ir {
+
+Load::Load(const type::Type* type, Value* f) : Base(), result_type(type), from(f) {
+    TINT_ASSERT(IR, result_type);
+    TINT_ASSERT(IR, from);
+    from->AddUsage(this);
 }
 
-}  // namespace
-}  // namespace tint::writer::spirv
+Load::~Load() = default;
+
+}  // namespace tint::ir
diff --git a/src/tint/ir/load.h b/src/tint/ir/load.h
new file mode 100644
index 0000000..b15eced
--- /dev/null
+++ b/src/tint/ir/load.h
@@ -0,0 +1,49 @@
+// Copyright 2023 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_TINT_IR_LOAD_H_
+#define SRC_TINT_IR_LOAD_H_
+
+#include "src/tint/ir/instruction.h"
+#include "src/tint/utils/castable.h"
+
+namespace tint::ir {
+
+/// A load instruction in the IR.
+class Load : public utils::Castable<Load, Instruction> {
+  public:
+    /// Constructor
+    /// @param type the result type
+    /// @param from the value being loaded from
+    Load(const type::Type* type, Value* from);
+    Load(const Load& inst) = delete;
+    Load(Load&& inst) = delete;
+    ~Load() override;
+
+    Load& operator=(const Load& inst) = delete;
+    Load& operator=(Load&& inst) = delete;
+
+    /// @returns the type of the value
+    const type::Type* Type() const override { return result_type; }
+
+    /// the result type of the instruction
+    const type::Type* result_type = nullptr;
+
+    /// the value being loaded
+    Value* from = nullptr;
+};
+
+}  // namespace tint::ir
+
+#endif  // SRC_TINT_IR_LOAD_H_
diff --git a/src/tint/ir/load_test.cc b/src/tint/ir/load_test.cc
new file mode 100644
index 0000000..2c6e5c1
--- /dev/null
+++ b/src/tint/ir/load_test.cc
@@ -0,0 +1,59 @@
+// Copyright 2023 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/tint/ir/builder.h"
+#include "src/tint/ir/instruction.h"
+#include "src/tint/ir/test_helper.h"
+
+namespace tint::ir {
+namespace {
+
+using namespace tint::number_suffixes;  // NOLINT
+
+using IR_InstructionTest = TestHelper;
+
+TEST_F(IR_InstructionTest, CreateLoad) {
+    Module mod;
+    Builder b{mod};
+
+    auto* store_type = b.ir.types.Get<type::I32>();
+    auto* var = b.Declare(b.ir.types.Get<type::Pointer>(
+        store_type, builtin::AddressSpace::kFunction, builtin::Access::kReadWrite));
+    const auto* inst = b.Load(var);
+
+    ASSERT_TRUE(inst->Is<Load>());
+    ASSERT_EQ(inst->from, var);
+
+    EXPECT_EQ(inst->Type(), store_type);
+
+    ASSERT_TRUE(inst->from->Is<ir::Var>());
+    EXPECT_EQ(inst->from, var);
+}
+
+TEST_F(IR_InstructionTest, Load_Usage) {
+    Module mod;
+    Builder b{mod};
+
+    auto* store_type = b.ir.types.Get<type::I32>();
+    auto* var = b.Declare(b.ir.types.Get<type::Pointer>(
+        store_type, builtin::AddressSpace::kFunction, builtin::Access::kReadWrite));
+    const auto* inst = b.Load(var);
+
+    ASSERT_NE(inst->from, nullptr);
+    ASSERT_EQ(inst->from->Usage().Length(), 1u);
+    EXPECT_EQ(inst->from->Usage()[0], inst);
+}
+
+}  // namespace
+}  // namespace tint::ir
diff --git a/src/tint/ir/module.h b/src/tint/ir/module.h
index a7929d2..244bc0a 100644
--- a/src/tint/ir/module.h
+++ b/src/tint/ir/module.h
@@ -70,8 +70,6 @@
     utils::BlockAllocator<constant::Value> constants;
     /// The value allocator
     utils::BlockAllocator<Value> values;
-    /// The instruction allocator
-    utils::BlockAllocator<Instruction> instructions;
 
     /// List of functions in the program
     utils::Vector<Function*, 8> functions;
diff --git a/src/tint/ir/module_test.cc b/src/tint/ir/module_test.cc
index c9b1150..9cd36ae 100644
--- a/src/tint/ir/module_test.cc
+++ b/src/tint/ir/module_test.cc
@@ -25,23 +25,20 @@
 
 TEST_F(IR_ModuleTest, NameOfUnnamed) {
     Module mod;
-    auto* v = mod.values.Create<ir::Var>(
-        mod.types.Get<type::I32>(), builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
+    auto* v = mod.values.Create<ir::Var>(mod.types.Get<type::I32>());
     EXPECT_FALSE(mod.NameOf(v).IsValid());
 }
 
 TEST_F(IR_ModuleTest, SetName) {
     Module mod;
-    auto* v = mod.values.Create<ir::Var>(
-        mod.types.Get<type::I32>(), builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
+    auto* v = mod.values.Create<ir::Var>(mod.types.Get<type::I32>());
     EXPECT_EQ(mod.SetName(v, "a").Name(), "a");
     EXPECT_EQ(mod.NameOf(v).Name(), "a");
 }
 
 TEST_F(IR_ModuleTest, SetNameRename) {
     Module mod;
-    auto* v = mod.values.Create<ir::Var>(
-        mod.types.Get<type::I32>(), builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
+    auto* v = mod.values.Create<ir::Var>(mod.types.Get<type::I32>());
     EXPECT_EQ(mod.SetName(v, "a").Name(), "a");
     EXPECT_EQ(mod.SetName(v, "b").Name(), "b");
     EXPECT_EQ(mod.NameOf(v).Name(), "b");
@@ -49,12 +46,9 @@
 
 TEST_F(IR_ModuleTest, SetNameCollision) {
     Module mod;
-    auto* a = mod.values.Create<ir::Var>(
-        mod.types.Get<type::I32>(), builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
-    auto* b = mod.values.Create<ir::Var>(
-        mod.types.Get<type::I32>(), builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
-    auto* c = mod.values.Create<ir::Var>(
-        mod.types.Get<type::I32>(), builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
+    auto* a = mod.values.Create<ir::Var>(mod.types.Get<type::I32>());
+    auto* b = mod.values.Create<ir::Var>(mod.types.Get<type::I32>());
+    auto* c = mod.values.Create<ir::Var>(mod.types.Get<type::I32>());
     EXPECT_EQ(mod.SetName(a, "x").Name(), "x");
     EXPECT_EQ(mod.SetName(b, "x_1").Name(), "x_1");
     EXPECT_EQ(mod.SetName(c, "x").Name(), "x_2");
diff --git a/src/tint/ir/to_program.cc b/src/tint/ir/to_program.cc
index 8a2dca9..b1aa67f 100644
--- a/src/tint/ir/to_program.cc
+++ b/src/tint/ir/to_program.cc
@@ -14,15 +14,19 @@
 
 #include "src/tint/ir/to_program.h"
 
+#include <string>
 #include <utility>
 
 #include "src/tint/ir/block.h"
 #include "src/tint/ir/call.h"
 #include "src/tint/ir/constant.h"
+#include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/instruction.h"
+#include "src/tint/ir/load.h"
 #include "src/tint/ir/module.h"
 #include "src/tint/ir/store.h"
+#include "src/tint/ir/switch.h"
 #include "src/tint/ir/user_call.h"
 #include "src/tint/ir/var.h"
 #include "src/tint/program_builder.h"
@@ -40,6 +44,17 @@
 #include "src/tint/utils/transform.h"
 #include "src/tint/utils/vector.h"
 
+// Helper for calling TINT_UNIMPLEMENTED() from a Switch(object_ptr) default case.
+#define UNHANDLED_CASE(object_ptr)          \
+    TINT_UNIMPLEMENTED(IR, b.Diagnostics()) \
+        << "unhandled case in Switch(): " << (object_ptr ? object_ptr->TypeInfo().name : "<null>")
+
+// Helper for incrementing nesting_depth_ and then decrementing nesting_depth_ at the end
+// of the scope that holds the call.
+#define SCOPED_NESTING() \
+    nesting_depth_++;    \
+    TINT_DEFER(nesting_depth_--)
+
 namespace tint::ir {
 
 namespace {
@@ -58,76 +73,264 @@
     }
 
   private:
+    /// The source IR module
     const Module& mod;
+
+    /// The target ProgramBuilder
     ProgramBuilder b;
+
+    /// A hashmap of value to symbol used in the emitted AST
     utils::Hashmap<const Value*, Symbol, 32> value_names_;
 
-    void Fn(const Function* fn) {
+    // The nesting depth of the currently generated AST
+    // 0 is module scope
+    // 1 is root-level function scope
+    // 2+ is within control flow
+    uint32_t nesting_depth_ = 0;
+
+    const ast::Function* Fn(const Function* fn) {
+        SCOPED_NESTING();
+
         auto name = Sym(fn->name);
         // TODO(crbug.com/tint/1915): Properly implement this when we've fleshed out Function
         utils::Vector<const ast::Parameter*, 1> params{};
-        ast::Type ret_ty;
-        auto* body = Block(fn->start_target);
+        auto ret_ty = Type(fn->return_type);
+        if (!ret_ty) {
+            return nullptr;
+        }
+        auto* body = FlowNodeGraph(fn->start_target);
+        if (!body) {
+            return nullptr;
+        }
         utils::Vector<const ast::Attribute*, 1> attrs{};
         utils::Vector<const ast::Attribute*, 1> ret_attrs{};
-        b.Func(name, std::move(params), ret_ty, body, std::move(attrs), std::move(ret_attrs));
+        return b.Func(name, std::move(params), ret_ty.Get(), body, std::move(attrs),
+                      std::move(ret_attrs));
     }
 
-    const ast::BlockStatement* Block(const ir::Block* block) {
+    const ast::BlockStatement* FlowNodeGraph(ir::FlowNode* start_node,
+                                             ir::FlowNode* stop_at = nullptr) {
         // TODO(crbug.com/tint/1902): Check if the block is dead
-        utils::Vector<const ast::Statement*, decltype(ir::Block::instructions)::static_length>
+        utils::Vector<const ast::Statement*,
+                      decltype(ast::BlockStatement::statements)::static_length>
             stmts;
-        for (auto* inst : block->instructions) {
-            auto* stmt = Stmt(inst);
-            if (!stmt) {
+
+        ir::Branch root_branch{start_node, {}};
+        const ir::Branch* branch = &root_branch;
+
+        // TODO(crbug.com/tint/1902): Handle block arguments.
+
+        while (branch->target != stop_at) {
+            enum Status { kContinue, kStop, kError };
+            Status status = tint::Switch(
+                branch->target,
+
+                [&](const ir::Block* block) {
+                    for (auto* inst : block->instructions) {
+                        auto stmt = Stmt(inst);
+                        if (TINT_UNLIKELY(!stmt)) {
+                            return kError;
+                        }
+                        if (auto* s = stmt.Get()) {
+                            stmts.Push(s);
+                        }
+                    }
+                    branch = &block->branch;
+                    return kContinue;
+                },
+
+                [&](const ir::If* if_) {
+                    auto* stmt = If(if_);
+                    if (TINT_UNLIKELY(!stmt)) {
+                        return kError;
+                    }
+                    stmts.Push(stmt);
+                    branch = &if_->merge;
+                    return branch->target->inbound_branches.IsEmpty() ? kStop : kContinue;
+                },
+
+                [&](const ir::Switch* switch_) {
+                    auto* stmt = Switch(switch_);
+                    if (TINT_UNLIKELY(!stmt)) {
+                        return kError;
+                    }
+                    stmts.Push(stmt);
+                    branch = &switch_->merge;
+                    return branch->target->inbound_branches.IsEmpty() ? kStop : kContinue;
+                },
+
+                [&](const ir::FunctionTerminator*) {
+                    auto res = FunctionTerminator(branch);
+                    if (TINT_UNLIKELY(!res)) {
+                        return kError;
+                    }
+                    if (auto* stmt = res.Get()) {
+                        stmts.Push(stmt);
+                    }
+                    return kStop;
+                },
+
+                [&](Default) {
+                    UNHANDLED_CASE(branch->target);
+                    return kError;
+                });
+
+            if (TINT_UNLIKELY(status == kError)) {
                 return nullptr;
             }
-            stmts.Push(stmt);
+            if (status == kStop) {
+                break;
+            }
         }
+
         return b.Block(std::move(stmts));
     }
 
-    const ast::Statement* FlowNode(const ir::FlowNode* node) {
-        // TODO(crbug.com/tint/1902): Check the node is connected
-        return Switch(
-            node,  //
-            [&](const ir::If* i) {
-                auto* cond = Expr(i->condition);
-                auto* t = Branch(i->true_);
-                if (auto* f = Branch(i->false_)) {
-                    return b.If(cond, t, b.Else(f));
-                }
-                // TODO(crbug.com/tint/1902): Emit merge block
-                return b.If(cond, t);
-            },
-            [&](Default) {
-                TINT_UNIMPLEMENTED(IR, b.Diagnostics())
-                    << "unhandled case in Switch(): " << node->TypeInfo().name;
-                return nullptr;
-            });
-    }
+    const ast::IfStatement* If(const ir::If* i) {
+        SCOPED_NESTING();
 
-    const ast::BlockStatement* Branch(const ir::Branch& branch) {
-        auto* stmt = FlowNode(branch.target);
-        if (!stmt) {
+        auto* cond = Expr(i->condition);
+        auto* t = FlowNodeGraph(i->true_.target, i->merge.target);
+        if (TINT_UNLIKELY(!t)) {
             return nullptr;
         }
-        if (auto* block = stmt->As<ast::BlockStatement>()) {
-            return block;
+
+        if (!IsEmpty(i->false_.target, i->merge.target)) {
+            // If the else target is an if flow node with the same merge target as this if, then
+            // emit an 'else if' instead of a block statement for the else.
+            if (auto* else_if = As<ir::If>(NextNonEmptyNode(i->false_.target));
+                else_if &&
+                NextNonEmptyNode(i->merge.target) == NextNonEmptyNode(else_if->merge.target)) {
+                auto* f = If(else_if);
+                if (!f) {
+                    return nullptr;
+                }
+                return b.If(cond, t, b.Else(f));
+            } else {
+                auto* f = FlowNodeGraph(i->false_.target, i->merge.target);
+                if (!f) {
+                    return nullptr;
+                }
+                return b.If(cond, t, b.Else(f));
+            }
         }
-        return b.Block(stmt);
+
+        return b.If(cond, t);
     }
 
-    const ast::Statement* Stmt(const ir::Instruction* inst) {
-        return Switch(
+    const ast::SwitchStatement* Switch(const ir::Switch* s) {
+        SCOPED_NESTING();
+
+        auto* cond = Expr(s->condition);
+        if (!cond) {
+            return nullptr;
+        }
+
+        auto cases = utils::Transform(
+            s->cases,  //
+            [&](const ir::Switch::Case& c) -> const tint::ast::CaseStatement* {
+                SCOPED_NESTING();
+                auto* body = FlowNodeGraph(c.start.target, s->merge.target);
+                if (!body) {
+                    return nullptr;
+                }
+
+                auto selectors = utils::Transform(
+                    c.selectors,  //
+                    [&](const ir::Switch::CaseSelector& cs) -> const ast::CaseSelector* {
+                        if (cs.IsDefault()) {
+                            return b.DefaultCaseSelector();
+                        }
+                        auto* expr = Expr(cs.val);
+                        if (!expr) {
+                            return nullptr;
+                        }
+                        return b.CaseSelector(expr);
+                    });
+                if (selectors.Any(utils::IsNull)) {
+                    return nullptr;
+                }
+
+                return b.Case(std::move(selectors), body);
+            });
+        if (cases.Any(utils::IsNull)) {
+            return nullptr;
+        }
+
+        return b.Switch(cond, std::move(cases));
+    }
+
+    utils::Result<const ast::ReturnStatement*> FunctionTerminator(const ir::Branch* branch) {
+        if (branch->args.IsEmpty()) {
+            // Branch to function terminator has no arguments.
+            // If this block is nested withing some control flow, then we must emit a
+            // 'return' statement, otherwise we've just naturally reached the end of the
+            // function where the 'return' is redundant.
+            if (nesting_depth_ > 1) {
+                return b.Return();
+            }
+            return nullptr;
+        }
+
+        // Branch to function terminator has arguments - this is the return value.
+        if (branch->args.Length() != 1) {
+            TINT_ICE(IR, b.Diagnostics())
+                << "expected 1 value for function terminator (return value), got "
+                << branch->args.Length();
+            return utils::Failure;
+        }
+
+        auto* val = Expr(branch->args.Front());
+        if (TINT_UNLIKELY(!val)) {
+            return utils::Failure;
+        }
+
+        return b.Return(val);
+    }
+
+    /// @return true if there are no instructions between @p node and and @p stop_at
+    bool IsEmpty(const ir::FlowNode* node, const ir::FlowNode* stop_at) {
+        while (node != stop_at) {
+            if (auto* block = node->As<ir::Block>()) {
+                if (block->instructions.Length() > 0) {
+                    return false;
+                }
+                node = block->branch.target;
+            } else {
+                return false;
+            }
+        }
+        return true;
+    }
+
+    /// @return the next flow node that isn't an empty block
+    const ir::FlowNode* NextNonEmptyNode(const ir::FlowNode* node) {
+        while (node) {
+            if (auto* block = node->As<ir::Block>()) {
+                for (auto* inst : block->instructions) {
+                    // Load instructions will be inlined, so ignore them.
+                    if (!inst->Is<ir::Load>()) {
+                        return node;
+                    }
+                }
+                node = block->branch.target;
+            } else {
+                return node;
+            }
+        }
+        return nullptr;
+    }
+
+    utils::Result<const ast::Statement*> Stmt(const ir::Instruction* inst) {
+        return tint::Switch<utils::Result<const ast::Statement*>>(
             inst,                                            //
             [&](const ir::Call* i) { return CallStmt(i); },  //
             [&](const ir::Var* i) { return Var(i); },        //
-            [&](const ir::Store* i) { return Store(i); },
+            [&](const ir::Load*) { return nullptr; },
+            [&](const ir::Store* i) { return Store(i); },  //
             [&](Default) {
-                TINT_UNIMPLEMENTED(IR, b.Diagnostics())
-                    << "unhandled case in Switch(): " << inst->TypeInfo().name;
-                return nullptr;
+                UNHANDLED_CASE(inst);
+                return utils::Failure;
             });
     }
 
@@ -141,7 +344,12 @@
 
     const ast::VariableDeclStatement* Var(const ir::Var* var) {
         Symbol name = NameOf(var);
-        auto ty = Type(var->Type());
+        auto* ptr = var->Type()->As<type::Pointer>();
+        if (!ptr) {
+            Err("Incorrect type for var");
+            return nullptr;
+        }
+        auto ty = Type(ptr->StoreType());
         const ast::Expression* init = nullptr;
         if (var->initializer) {
             init = Expr(var->initializer);
@@ -149,13 +357,13 @@
                 return nullptr;
             }
         }
-        switch (var->address_space) {
+        switch (ptr->AddressSpace()) {
             case builtin::AddressSpace::kFunction:
-                return b.Decl(b.Var(name, ty, init));
+                return b.Decl(b.Var(name, ty.Get(), init));
             case builtin::AddressSpace::kStorage:
-                return b.Decl(b.Var(name, ty, init, var->access, var->address_space));
+                return b.Decl(b.Var(name, ty.Get(), init, ptr->Access(), ptr->AddressSpace()));
             default:
-                return b.Decl(b.Var(name, ty, init, var->address_space));
+                return b.Decl(b.Var(name, ty.Get(), init, ptr->AddressSpace()));
         }
     }
 
@@ -169,29 +377,29 @@
         if (args.Any(utils::IsNull)) {
             return nullptr;
         }
-        return Switch(
+        return tint::Switch(
             call,  //
             [&](const ir::UserCall* c) { return b.Call(Sym(c->name), std::move(args)); },
             [&](Default) {
-                TINT_UNIMPLEMENTED(IR, b.Diagnostics())
-                    << "unhandled case in Switch(): " << call->TypeInfo().name;
+                UNHANDLED_CASE(call);
                 return nullptr;
             });
     }
 
     const ast::Expression* Expr(const ir::Value* val) {
-        return Switch(
+        return tint::Switch(
             val,  //
             [&](const ir::Constant* c) { return ConstExpr(c); },
+            [&](const ir::Load* l) { return LoadExpr(l); },
+            [&](const ir::Var* v) { return VarExpr(v); },
             [&](Default) {
-                TINT_UNIMPLEMENTED(IR, b.Diagnostics())
-                    << "unhandled case in Switch(): " << val->TypeInfo().name;
+                UNHANDLED_CASE(val);
                 return nullptr;
             });
     }
 
     const ast::Expression* ConstExpr(const ir::Constant* c) {
-        return Switch(
+        return tint::Switch(
             c->Type(),  //
             [&](const type::I32*) { return b.Expr(c->value->ValueAs<i32>()); },
             [&](const type::U32*) { return b.Expr(c->value->ValueAs<u32>()); },
@@ -199,14 +407,17 @@
             [&](const type::F16*) { return b.Expr(c->value->ValueAs<f16>()); },
             [&](const type::Bool*) { return b.Expr(c->value->ValueAs<bool>()); },
             [&](Default) {
-                TINT_UNIMPLEMENTED(IR, b.Diagnostics())
-                    << "unhandled case in Switch(): " << c->TypeInfo().name;
+                UNHANDLED_CASE(c);
                 return nullptr;
             });
     }
 
-    const ast::Type Type(const type::Type* ty) {
-        return Switch(
+    const ast::Expression* LoadExpr(const ir::Load* l) { return Expr(l->from); }
+
+    const ast::Expression* VarExpr(const ir::Var* v) { return b.Expr(NameOf(v)); }
+
+    utils::Result<ast::Type> Type(const type::Type* ty) {
+        return tint::Switch<utils::Result<ast::Type>>(
             ty,                                              //
             [&](const type::Void*) { return ast::Type{}; },  //
             [&](const type::I32*) { return b.ty.i32(); },    //
@@ -214,64 +425,94 @@
             [&](const type::F16*) { return b.ty.f16(); },    //
             [&](const type::F32*) { return b.ty.f32(); },    //
             [&](const type::Bool*) { return b.ty.bool_(); },
-            [&](const type::Matrix* m) {
+            [&](const type::Matrix* m) -> utils::Result<ast::Type> {
                 auto el = Type(m->type());
-                return b.ty.mat(el, m->columns(), m->rows());
+                if (!el) {
+                    return utils::Failure;
+                }
+                return b.ty.mat(el.Get(), m->columns(), m->rows());
             },
-            [&](const type::Vector* v) {
+            [&](const type::Vector* v) -> utils::Result<ast::Type> {
                 auto el = Type(v->type());
+                if (!el) {
+                    return utils::Failure;
+                }
                 if (v->Packed()) {
                     TINT_ASSERT(IR, v->Width() == 3u);
-                    return b.ty(builtin::Builtin::kPackedVec3, el);
+                    return b.ty(builtin::Builtin::kPackedVec3, el.Get());
                 } else {
-                    return b.ty.vec(el, v->Width());
+                    return b.ty.vec(el.Get(), v->Width());
                 }
             },
-            [&](const type::Array* a) {
+            [&](const type::Array* a) -> utils::Result<ast::Type> {
                 auto el = Type(a->ElemType());
+                if (!el) {
+                    return utils::Failure;
+                }
                 utils::Vector<const ast::Attribute*, 1> attrs;
                 if (!a->IsStrideImplicit()) {
                     attrs.Push(b.Stride(a->Stride()));
                 }
                 if (a->Count()->Is<type::RuntimeArrayCount>()) {
-                    return b.ty.array(el, std::move(attrs));
+                    return b.ty.array(el.Get(), std::move(attrs));
                 }
                 auto count = a->ConstantCount();
                 if (TINT_UNLIKELY(!count)) {
                     TINT_ICE(IR, b.Diagnostics()) << type::Array::kErrExpectedConstantCount;
-                    return b.ty.array(el, u32(1), std::move(attrs));
+                    return b.ty.array(el.Get(), u32(1), std::move(attrs));
                 }
-                return b.ty.array(el, u32(count.value()), std::move(attrs));
+                return b.ty.array(el.Get(), u32(count.value()), std::move(attrs));
             },
             [&](const type::Struct* s) { return b.ty(s->Name().NameView()); },
-            [&](const type::Atomic* a) { return b.ty.atomic(Type(a->Type())); },
+            [&](const type::Atomic* a) -> utils::Result<ast::Type> {
+                auto el = Type(a->Type());
+                if (!el) {
+                    return utils::Failure;
+                }
+                return b.ty.atomic(el.Get());
+            },
             [&](const type::DepthTexture* t) { return b.ty.depth_texture(t->dim()); },
             [&](const type::DepthMultisampledTexture* t) {
                 return b.ty.depth_multisampled_texture(t->dim());
             },
             [&](const type::ExternalTexture*) { return b.ty.external_texture(); },
-            [&](const type::MultisampledTexture* t) {
-                return b.ty.multisampled_texture(t->dim(), Type(t->type()));
+            [&](const type::MultisampledTexture* t) -> utils::Result<ast::Type> {
+                auto el = Type(t->type());
+                if (!el) {
+                    return utils::Failure;
+                }
+                return b.ty.multisampled_texture(t->dim(), el.Get());
             },
-            [&](const type::SampledTexture* t) {
-                return b.ty.sampled_texture(t->dim(), Type(t->type()));
+            [&](const type::SampledTexture* t) -> utils::Result<ast::Type> {
+                auto el = Type(t->type());
+                if (!el) {
+                    return utils::Failure;
+                }
+                return b.ty.sampled_texture(t->dim(), el.Get());
             },
             [&](const type::StorageTexture* t) {
                 return b.ty.storage_texture(t->dim(), t->texel_format(), t->access());
             },
             [&](const type::Sampler* s) { return b.ty.sampler(s->kind()); },
-            [&](const type::Pointer* p) {
+            [&](const type::Pointer* p) -> utils::Result<ast::Type> {
                 // Note: type::Pointer always has an inferred access, but WGSL only allows an
                 // explicit access in the 'storage' address space.
+                auto el = Type(p->StoreType());
+                if (!el) {
+                    return utils::Failure;
+                }
                 auto address_space = p->AddressSpace();
                 auto access = address_space == builtin::AddressSpace::kStorage
                                   ? p->Access()
                                   : builtin::Access::kUndefined;
-                return b.ty.pointer(Type(p->StoreType()), address_space, access);
+                return b.ty.pointer(el.Get(), address_space, access);
             },
-            [&](const type::Reference* r) { return Type(r->StoreType()); },
+            [&](const type::Reference*) -> utils::Result<ast::Type> {
+                TINT_ICE(IR, b.Diagnostics()) << "reference types should never appear in the IR";
+                return ast::Type{};
+            },
             [&](Default) {
-                TINT_UNREACHABLE(IR, b.Diagnostics()) << "unhandled type: " << ty->TypeInfo().name;
+                UNHANDLED_CASE(ty);
                 return ast::Type{};
             });
     }
@@ -288,7 +529,7 @@
 
     Symbol Sym(const Symbol& s) { return b.Symbols().Register(s.NameView()); }
 
-    // void Err(std::string str) { b.Diagnostics().add_error(diag::System::IR, std::move(str)); }
+    void Err(std::string str) { b.Diagnostics().add_error(diag::System::IR, std::move(str)); }
 };
 
 }  // namespace
diff --git a/src/tint/ir/to_program_roundtrip_test.cc b/src/tint/ir/to_program_roundtrip_test.cc
index 11aff46..8f26792 100644
--- a/src/tint/ir/to_program_roundtrip_test.cc
+++ b/src/tint/ir/to_program_roundtrip_test.cc
@@ -40,6 +40,13 @@
         ASSERT_TRUE(ir_module);
 
         auto output_program = ToProgram(ir_module.Get());
+        if (!output_program.IsValid()) {
+            tint::ir::Disassembler d{ir_module.Get()};
+            FAIL() << output_program.Diagnostics().str() << std::endl
+                   << "IR:" << std::endl
+                   << d.Disassemble();
+        }
+
         ASSERT_TRUE(output_program.IsValid()) << output_program.Diagnostics().str();
 
         auto output = writer::wgsl::Generate(&output_program, {});
@@ -60,13 +67,44 @@
     Test("");
 }
 
-TEST_F(IRToProgramRoundtripTest, EmptySingleFunction) {
+TEST_F(IRToProgramRoundtripTest, SingleFunction_Empty) {
     Test(R"(
 fn f() {
 }
 )");
 }
 
+TEST_F(IRToProgramRoundtripTest, SingleFunction_Return) {
+    Test(R"(
+fn f() {
+  return;
+}
+)",
+         R"(
+fn f() {
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, SingleFunction_Return_i32) {
+    Test(R"(
+fn f() -> i32 {
+  return 42i;
+}
+)");
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Function-scope var
+////////////////////////////////////////////////////////////////////////////////
+TEST_F(IRToProgramRoundtripTest, FunctionScopeVar_i32) {
+    Test(R"(
+fn f() {
+  var i : i32;
+}
+)");
+}
+
 TEST_F(IRToProgramRoundtripTest, FunctionScopeVar_i32_InitLiteral) {
     Test(R"(
 fn f() {
@@ -75,5 +113,256 @@
 )");
 }
 
+TEST_F(IRToProgramRoundtripTest, FunctionScopeVar_Chained) {
+    Test(R"(
+fn f() {
+  var a : i32 = 42i;
+  var b : i32 = a;
+  var c : i32 = b;
+}
+)");
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// If
+////////////////////////////////////////////////////////////////////////////////
+TEST_F(IRToProgramRoundtripTest, If_CallFn) {
+    Test(R"(
+fn a() {
+}
+
+fn f() {
+  var cond : bool = true;
+  if (cond) {
+    a();
+  }
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, If_Return) {
+    Test(R"(
+fn f() {
+  var cond : bool = true;
+  if (cond) {
+    return;
+  }
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, If_Return_i32) {
+    Test(R"(
+fn f() -> i32 {
+  var cond : bool = true;
+  if (cond) {
+    return 42i;
+  }
+  return 10i;
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, If_CallFn_Else_CallFn) {
+    Test(R"(
+fn a() {
+}
+
+fn b() {
+}
+
+fn f() {
+  var cond : bool = true;
+  if (cond) {
+    a();
+  } else {
+    b();
+  }
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, If_Return_f32_Else_Return_f32) {
+    Test(R"(
+fn f() -> f32 {
+  var cond : bool = true;
+  if (cond) {
+    return 1.0f;
+  } else {
+    return 2.0f;
+  }
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, If_Return_u32_Else_CallFn) {
+    Test(R"(
+fn a() {
+}
+
+fn b() {
+}
+
+fn f() -> u32 {
+  var cond : bool = true;
+  if (cond) {
+    return 1u;
+  } else {
+    a();
+  }
+  b();
+  return 2u;
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, If_CallFn_ElseIf_CallFn) {
+    Test(R"(
+fn a() {
+}
+
+fn b() {
+}
+
+fn c() {
+}
+
+fn f() {
+  var cond_a : bool = true;
+  var cond_b : bool = true;
+  if (cond_a) {
+    a();
+  } else if (cond_b) {
+    b();
+  }
+  c();
+}
+)");
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Switch
+////////////////////////////////////////////////////////////////////////////////
+TEST_F(IRToProgramRoundtripTest, Switch_Default) {
+    Test(R"(
+fn a() {
+}
+
+fn f() {
+  var v : i32 = 42i;
+  switch(v) {
+    default: {
+      a();
+    }
+  }
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, Switch_3_Cases) {
+    Test(R"(
+fn a() {
+}
+
+fn b() {
+}
+
+fn c() {
+}
+
+fn f() {
+  var v : i32 = 42i;
+  switch(v) {
+    case 0i: {
+      a();
+    }
+    case 1i, default: {
+      b();
+    }
+    case 2i: {
+      c();
+    }
+  }
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, Switch_3_Cases_AllReturn) {
+    Test(R"(
+fn a() {
+}
+
+fn f() {
+  var v : i32 = 42i;
+  switch(v) {
+    case 0i: {
+      return;
+    }
+    case 1i, default: {
+      return;
+    }
+    case 2i: {
+      return;
+    }
+  }
+  a();
+}
+)",
+         R"(
+fn a() {
+}
+
+fn f() {
+  var v : i32 = 42i;
+  switch(v) {
+    case 0i: {
+      return;
+    }
+    case 1i, default: {
+      return;
+    }
+    case 2i: {
+      return;
+    }
+  }
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, Switch_Nested) {
+    Test(R"(
+fn a() {
+}
+
+fn b() {
+}
+
+fn c() {
+}
+
+fn f() {
+  var v1 : i32 = 42i;
+  var v2 : i32 = 24i;
+  switch(v1) {
+    case 0i: {
+      a();
+    }
+    case 1i, default: {
+      switch(v2) {
+        case 0i: {
+        }
+        case 1i, default: {
+          return;
+        }
+      }
+    }
+    case 2i: {
+      c();
+    }
+  }
+}
+)");
+}
+
 }  // namespace
 }  // namespace tint::ir
diff --git a/src/tint/ir/transform/add_empty_entry_point.cc b/src/tint/ir/transform/add_empty_entry_point.cc
new file mode 100644
index 0000000..809a6ad
--- /dev/null
+++ b/src/tint/ir/transform/add_empty_entry_point.cc
@@ -0,0 +1,45 @@
+// Copyright 2023 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/tint/ir/transform/add_empty_entry_point.h"
+
+#include <utility>
+
+#include "src/tint/ir/builder.h"
+#include "src/tint/ir/module.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::transform::AddEmptyEntryPoint);
+
+namespace tint::ir::transform {
+
+AddEmptyEntryPoint::AddEmptyEntryPoint() = default;
+
+AddEmptyEntryPoint::~AddEmptyEntryPoint() = default;
+
+void AddEmptyEntryPoint::Run(ir::Module* ir, const DataMap&, DataMap&) const {
+    for (auto* func : ir->functions) {
+        if (func->pipeline_stage != Function::PipelineStage::kUndefined) {
+            return;
+        }
+    }
+
+    ir::Builder builder(*ir);
+    auto* ep =
+        builder.CreateFunction(ir->symbols.New("unused_entry_point"), ir->types.Get<type::Void>(),
+                               Function::PipelineStage::kCompute, std::array{1u, 1u, 1u});
+    builder.Branch(ep->start_target, ep->end_target);
+    ir->functions.Push(ep);
+}
+
+}  // namespace tint::ir::transform
diff --git a/src/tint/ir/transform/add_empty_entry_point.h b/src/tint/ir/transform/add_empty_entry_point.h
new file mode 100644
index 0000000..39f3413
--- /dev/null
+++ b/src/tint/ir/transform/add_empty_entry_point.h
@@ -0,0 +1,36 @@
+// Copyright 2023 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_TINT_IR_TRANSFORM_ADD_EMPTY_ENTRY_POINT_H_
+#define SRC_TINT_IR_TRANSFORM_ADD_EMPTY_ENTRY_POINT_H_
+
+#include "src/tint/ir/transform/transform.h"
+
+namespace tint::ir::transform {
+
+/// Add an empty entry point to the module, if no other entry points exist.
+class AddEmptyEntryPoint final : public utils::Castable<AddEmptyEntryPoint, Transform> {
+  public:
+    /// Constructor
+    AddEmptyEntryPoint();
+    /// Destructor
+    ~AddEmptyEntryPoint() override;
+
+    /// @copydoc Transform::Run
+    void Run(ir::Module* module, const DataMap& inputs, DataMap& outputs) const override;
+};
+
+}  // namespace tint::ir::transform
+
+#endif  // SRC_TINT_IR_TRANSFORM_ADD_EMPTY_ENTRY_POINT_H_
diff --git a/src/tint/ir/transform/add_empty_entry_point_test.cc b/src/tint/ir/transform/add_empty_entry_point_test.cc
new file mode 100644
index 0000000..baba8e0
--- /dev/null
+++ b/src/tint/ir/transform/add_empty_entry_point_test.cc
@@ -0,0 +1,60 @@
+// Copyright 2023 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/tint/ir/transform/add_empty_entry_point.h"
+
+#include <utility>
+
+#include "src/tint/ir/transform/test_helper.h"
+
+namespace tint::ir::transform {
+namespace {
+
+using IR_AddEmptyEntryPointTest = TransformTest;
+
+TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) {
+    auto* expect = R"(
+%fn1 = func unused_entry_point():void [@compute @workgroup_size(1, 1, 1)] {
+  %fn2 = block {
+  } -> %func_end # return
+} %func_end
+
+)";
+
+    Run<AddEmptyEntryPoint>();
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_AddEmptyEntryPointTest, ExistingEntryPoint) {
+    auto* ep = b.CreateFunction(mod.symbols.New("main"), mod.types.Get<type::Void>(),
+                                Function::PipelineStage::kFragment);
+    b.Branch(ep->start_target, ep->end_target);
+    mod.functions.Push(ep);
+
+    auto* expect = R"(
+%fn1 = func main():void [@fragment] {
+  %fn2 = block {
+  } -> %func_end # return
+} %func_end
+
+)";
+
+    Run<AddEmptyEntryPoint>();
+
+    EXPECT_EQ(expect, str());
+}
+
+}  // namespace
+}  // namespace tint::ir::transform
diff --git a/src/tint/ir/transform/test_helper.h b/src/tint/ir/transform/test_helper.h
new file mode 100644
index 0000000..a122eb2
--- /dev/null
+++ b/src/tint/ir/transform/test_helper.h
@@ -0,0 +1,72 @@
+// Copyright 2023 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_TINT_IR_TRANSFORM_TEST_HELPER_H_
+#define SRC_TINT_IR_TRANSFORM_TEST_HELPER_H_
+
+#include <memory>
+#include <string>
+#include <utility>
+#include <vector>
+
+#include "gtest/gtest.h"
+#include "src/tint/ir/builder.h"
+#include "src/tint/ir/disassembler.h"
+#include "src/tint/ir/transform/transform.h"
+#include "src/tint/transform/manager.h"
+
+namespace tint::ir::transform {
+
+/// Helper class for testing IR transforms.
+template <typename BASE>
+class TransformTestBase : public BASE {
+  public:
+    /// Transforms the module, using transforms in `TRANSFORMS`.
+    /// @param data the optional Transform::DataMap to pass to Transform::Run()
+    /// @returns the transform outputs, if any
+    template <typename... TRANSFORMS>
+    Transform::DataMap Run(const Transform::DataMap& data = {}) {
+        tint::transform::Manager manager;
+        tint::transform::DataMap outputs;
+        for (auto* transform_ptr : std::initializer_list<Transform*>{new TRANSFORMS()...}) {
+            manager.append(std::unique_ptr<Transform>(transform_ptr));
+        }
+        manager.Run(&mod, data, outputs);
+        return outputs;
+    }
+
+    /// @returns the transformed module as a disassembled string
+    std::string str() {
+        ir::Disassembler dis(mod);
+        return "\n" + dis.Disassemble();
+    }
+
+  protected:
+    /// The test IR module.
+    ir::Module mod;
+    /// The test IR builder.
+    ir::Builder b{mod};
+
+  private:
+    std::vector<std::unique_ptr<Source::File>> files_;
+};
+
+using TransformTest = TransformTestBase<testing::Test>;
+
+template <typename T>
+using TransformTestWithParam = TransformTestBase<testing::TestWithParam<T>>;
+
+}  // namespace tint::ir::transform
+
+#endif  // SRC_TINT_IR_TRANSFORM_TEST_HELPER_H_
diff --git a/src/tint/ir/transform/transform.cc b/src/tint/ir/transform/transform.cc
new file mode 100644
index 0000000..eb5e68e
--- /dev/null
+++ b/src/tint/ir/transform/transform.cc
@@ -0,0 +1,24 @@
+// Copyright 2023 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/tint/ir/transform/transform.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::transform::Transform);
+
+namespace tint::ir::transform {
+
+Transform::Transform() = default;
+Transform::~Transform() = default;
+
+}  // namespace tint::ir::transform
diff --git a/src/tint/ir/transform/transform.h b/src/tint/ir/transform/transform.h
new file mode 100644
index 0000000..a09fe77
--- /dev/null
+++ b/src/tint/ir/transform/transform.h
@@ -0,0 +1,48 @@
+// Copyright 2023 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_TINT_IR_TRANSFORM_TRANSFORM_H_
+#define SRC_TINT_IR_TRANSFORM_TRANSFORM_H_
+
+#include "src/tint/transform/transform.h"
+
+#include <utility>
+
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class Module;
+}  // namespace tint::ir
+
+namespace tint::ir::transform {
+
+/// Interface for IR Module transforms.
+class Transform : public utils::Castable<Transform, tint::transform::Transform> {
+  public:
+    /// Constructor
+    Transform();
+    /// Destructor
+    ~Transform() override;
+
+    /// Run the transform on @p module
+    /// @param module the source module to transform
+    /// @param inputs optional extra transform-specific input data
+    /// @param outputs optional extra transform-specific output data
+    virtual void Run(ir::Module* module, const DataMap& inputs, DataMap& outputs) const = 0;
+};
+
+}  // namespace tint::ir::transform
+
+#endif  // SRC_TINT_IR_TRANSFORM_TRANSFORM_H_
diff --git a/src/tint/ir/unary.h b/src/tint/ir/unary.h
index e665386..46edd45 100644
--- a/src/tint/ir/unary.h
+++ b/src/tint/ir/unary.h
@@ -25,9 +25,7 @@
   public:
     /// The kind of instruction.
     enum class Kind {
-        kAddressOf,
         kComplement,
-        kIndirection,
         kNegation,
     };
 
@@ -50,7 +48,7 @@
     const Value* Val() const { return val_; }
 
     /// the kind of unary instruction
-    Kind kind = Kind::kAddressOf;
+    Kind kind = Kind::kNegation;
 
     /// the result type of the instruction
     const type::Type* result_type = nullptr;
diff --git a/src/tint/ir/unary_test.cc b/src/tint/ir/unary_test.cc
index 392a75e..280c7de 100644
--- a/src/tint/ir/unary_test.cc
+++ b/src/tint/ir/unary_test.cc
@@ -23,27 +23,6 @@
 
 using IR_InstructionTest = TestHelper;
 
-TEST_F(IR_InstructionTest, CreateAddressOf) {
-    Module mod;
-    Builder b{mod};
-
-    // TODO(dsinclair): This would be better as an identifier, but works for now.
-    const auto* inst = b.AddressOf(
-        b.ir.types.Get<type::Pointer>(b.ir.types.Get<type::I32>(), builtin::AddressSpace::kPrivate,
-                                      builtin::Access::kReadWrite),
-        b.Constant(4_i));
-
-    ASSERT_TRUE(inst->Is<Unary>());
-    EXPECT_EQ(inst->kind, Unary::Kind::kAddressOf);
-
-    ASSERT_NE(inst->Type(), nullptr);
-
-    ASSERT_TRUE(inst->Val()->Is<Constant>());
-    auto lhs = inst->Val()->As<Constant>()->value;
-    ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
-    EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
-}
-
 TEST_F(IR_InstructionTest, CreateComplement) {
     Module mod;
     Builder b{mod};
@@ -58,22 +37,6 @@
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
 
-TEST_F(IR_InstructionTest, CreateIndirection) {
-    Module mod;
-    Builder b{mod};
-
-    // TODO(dsinclair): This would be better as an identifier, but works for now.
-    const auto* inst = b.Indirection(b.ir.types.Get<type::I32>(), b.Constant(4_i));
-
-    ASSERT_TRUE(inst->Is<Unary>());
-    EXPECT_EQ(inst->kind, Unary::Kind::kIndirection);
-
-    ASSERT_TRUE(inst->Val()->Is<Constant>());
-    auto lhs = inst->Val()->As<Constant>()->value;
-    ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
-    EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
-}
-
 TEST_F(IR_InstructionTest, CreateNegation) {
     Module mod;
     Builder b{mod};
diff --git a/src/tint/ir/var.cc b/src/tint/ir/var.cc
index e33da54..9bf4b8d 100644
--- a/src/tint/ir/var.cc
+++ b/src/tint/ir/var.cc
@@ -19,8 +19,7 @@
 
 namespace tint::ir {
 
-Var::Var(const type::Type* ty, builtin::AddressSpace addr_space, builtin::Access acc)
-    : type(ty), address_space(addr_space), access(acc) {}
+Var::Var(const type::Type* ty) : type(ty) {}
 
 Var::~Var() = default;
 
diff --git a/src/tint/ir/var.h b/src/tint/ir/var.h
index 5a61104..c874a62 100644
--- a/src/tint/ir/var.h
+++ b/src/tint/ir/var.h
@@ -27,9 +27,7 @@
   public:
     /// Constructor
     /// @param type the type of the var
-    /// @param address_space the address space of the var
-    /// @param access the access mode of the var
-    Var(const type::Type* type, builtin::AddressSpace address_space, builtin::Access access);
+    explicit Var(const type::Type* type);
     Var(const Var& inst) = delete;
     Var(Var&& inst) = delete;
     ~Var() override;
@@ -43,12 +41,6 @@
     /// the result type of the instruction
     const type::Type* type = nullptr;
 
-    /// The variable address space
-    builtin::AddressSpace address_space = builtin::AddressSpace::kUndefined;
-
-    /// The variable access mode
-    builtin::Access access = builtin::Access::kUndefined;
-
     /// The optional initializer
     Value* initializer = nullptr;
 };
diff --git a/src/tint/resolver/materialize_test.cc b/src/tint/resolver/materialize_test.cc
index 696f279..cf67c93 100644
--- a/src/tint/resolver/materialize_test.cc
+++ b/src/tint/resolver/materialize_test.cc
@@ -211,6 +211,10 @@
 
     // abstract_expr[runtime-index]
     kRuntimeIndex,
+
+    // var a : target_type;
+    // a += abstract_expr;
+    kCompoundAssign,
 };
 
 static std::ostream& operator<<(std::ostream& o, Method m) {
@@ -247,6 +251,8 @@
             return o << "workgroup-size";
         case Method::kRuntimeIndex:
             return o << "runtime-index";
+        case Method::kCompoundAssign:
+            return o << "compound-assign";
     }
     return o << "<unknown>";
 }
@@ -387,10 +393,15 @@
                  utils::Vector{WorkgroupSize(target_expr(), abstract_expr, Expr(123_a)),
                                Stage(ast::PipelineStage::kCompute)});
             break;
-        case Method::kRuntimeIndex:
+        case Method::kRuntimeIndex: {
             auto* runtime_index = Var("runtime_index", Expr(1_i));
             WrapInFunction(runtime_index, IndexAccessor(abstract_expr, runtime_index));
             break;
+        }
+        case Method::kCompoundAssign:
+            WrapInFunction(Decl(Var("a", target_ty())),
+                           CompoundAssign("a", abstract_expr, ast::BinaryOp::kAdd));
+            break;
     }
 
     switch (expectation) {
@@ -421,6 +432,10 @@
                     expect = "error: no matching overload for operator + (" +
                              data.target_type_name + ", " + data.abstract_type_name + ")";
                     break;
+                case Method::kCompoundAssign:
+                    expect = "error: no matching overload for operator += (" +
+                             data.target_type_name + ", " + data.abstract_type_name + ")";
+                    break;
                 default:
                     expect = "error: cannot convert value of type '" + data.abstract_type_name +
                              "' to type '" + data.target_type_name + "'";
@@ -440,13 +455,13 @@
 /// Methods that support scalar materialization
 constexpr Method kScalarMethods[] = {
     Method::kLet,    Method::kVar,   Method::kAssign, Method::kFnArg,    Method::kBuiltinArg,
-    Method::kReturn, Method::kArray, Method::kStruct, Method::kBinaryOp,
+    Method::kReturn, Method::kArray, Method::kStruct, Method::kBinaryOp, Method::kCompoundAssign,
 };
 
 /// Methods that support vector materialization
 constexpr Method kVectorMethods[] = {
     Method::kLet,    Method::kVar,   Method::kAssign, Method::kFnArg,    Method::kBuiltinArg,
-    Method::kReturn, Method::kArray, Method::kStruct, Method::kBinaryOp,
+    Method::kReturn, Method::kArray, Method::kStruct, Method::kBinaryOp, Method::kCompoundAssign,
 };
 
 /// Methods that support matrix materialization
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index a65d825..09d8464 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -4479,7 +4479,7 @@
             return false;
         }
 
-        auto* rhs = Load(ValueExpression(stmt->rhs));
+        const auto* rhs = ValueExpression(stmt->rhs);
         if (!rhs) {
             return false;
         }
@@ -4491,12 +4491,19 @@
         auto* lhs_ty = lhs->Type()->UnwrapRef();
         auto* rhs_ty = rhs->Type()->UnwrapRef();
         auto stage = sem::EarliestStage(lhs->Stage(), rhs->Stage());
-        auto* ty =
-            intrinsic_table_->Lookup(stmt->op, lhs_ty, rhs_ty, stage, stmt->source, true).result;
-        if (!ty) {
+
+        auto op = intrinsic_table_->Lookup(stmt->op, lhs_ty, rhs_ty, stage, stmt->source, true);
+        if (!op.result) {
             return false;
         }
-        return validator_.Assignment(stmt, ty);
+
+        // Load or materialize the RHS if necessary.
+        rhs = Load(Materialize(rhs, op.rhs));
+        if (!rhs) {
+            return false;
+        }
+
+        return validator_.Assignment(stmt, op.result);
     });
 }
 
diff --git a/src/tint/transform/manager.cc b/src/tint/transform/manager.cc
index eff3a0c..92d0427 100644
--- a/src/tint/transform/manager.cc
+++ b/src/tint/transform/manager.cc
@@ -17,12 +17,24 @@
 #include "src/tint/ast/transform/transform.h"
 #include "src/tint/program_builder.h"
 
+#if TINT_BUILD_IR
+#include "src/tint/ir/from_program.h"
+#include "src/tint/ir/to_program.h"
+#include "src/tint/ir/transform/transform.h"
+#else
+// Declare an ir::Module class so that the transform target variant compiles.
+namespace ir {
+class Module;
+}
+#endif  // TINT_BUILD_IR
+
 /// If set to 1 then the transform::Manager will dump the WGSL of the program
 /// before and after each transform. Helpful for debugging bad output.
 #define TINT_PRINT_PROGRAM_FOR_EACH_TRANSFORM 0
 
 #if TINT_PRINT_PROGRAM_FOR_EACH_TRANSFORM
 #include <iostream>
+#include "src/tint/ir/disassembler.h"
 #define TINT_IF_PRINT_PROGRAM(x) x
 #else  // TINT_PRINT_PROGRAM_FOR_EACH_TRANSFORM
 #define TINT_IF_PRINT_PROGRAM(x)
@@ -33,62 +45,140 @@
 Manager::Manager() = default;
 Manager::~Manager() = default;
 
-Program Manager::Run(const Program* program,
-                     const transform::DataMap& inputs,
-                     transform::DataMap& outputs) const {
+template <typename OUTPUT, typename INPUT>
+OUTPUT Manager::RunTransforms(INPUT in,
+                              const transform::DataMap& inputs,
+                              transform::DataMap& outputs) const {
+    static_assert(std::is_same<INPUT, const Program*>() || std::is_same<INPUT, ir::Module*>());
+    static_assert(std::is_same<OUTPUT, Program>() || std::is_same<OUTPUT, ir::Module*>());
+
+    // The current transform target, which could be either AST or IR.
+    std::variant<const Program*, ir::Module*> target = in;
+    // A local AST program to hold the result of AST transforms.
+    Program ast_result;
+#if TINT_BUILD_IR
+    // A local IR module to hold the result of AST->IR conversions.
+    ir::Module ir_result;
+#endif
+
 #if TINT_PRINT_PROGRAM_FOR_EACH_TRANSFORM
-    auto print_program = [&](const char* msg, const Transform* transform) {
-        auto wgsl = Program::printer(program);
+    auto print_program = [&](const char* msg, const char* name) {
         std::cout << "=========================================================" << std::endl;
-        std::cout << "== " << msg << " " << transform->TypeInfo().name << ":" << std::endl;
+        std::cout << "== " << msg << " " << name << ":" << std::endl;
         std::cout << "=========================================================" << std::endl;
-        std::cout << wgsl << std::endl;
-        if (!program->IsValid()) {
-            std::cout << "-- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- --" << std::endl;
-            std::cout << program->Diagnostics().str() << std::endl;
+        if (std::holds_alternative<const Program*>(target)) {
+            auto* program = std::get<const Program*>(target);
+            auto wgsl = Program::printer(program);
+            std::cout << wgsl << std::endl;
+            if (!program->IsValid()) {
+                std::cout << "-- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- --"
+                          << std::endl;
+                std::cout << program->Diagnostics().str() << std::endl;
+            }
+        } else if (std::holds_alternative<ir::Module*>(target)) {
+#if TINT_BUILD_IR
+            ir::Disassembler dis(*std::get<ir::Module*>(target));
+            std::cout << dis.Disassemble();
+#endif  // TINT_BUILD_IR
         }
         std::cout << "=========================================================" << std::endl
                   << std::endl;
     };
 #endif
 
-    std::optional<Program> output;
+    // Helper functions to get the current program state as either an AST program or IR module,
+    // performing a conversion if necessary.
+    auto get_ast = [&]() {
+#if TINT_BUILD_IR
+        if (std::holds_alternative<ir::Module*>(target)) {
+            // Convert the IR module to an AST program.
+            ast_result = ir::ToProgram(*std::get<ir::Module*>(target));
+            target = &ast_result;
+        }
+#endif  // TINT_BUILD_IR
+        TINT_ASSERT(Transform, std::holds_alternative<const Program*>(target));
+        return std::get<const Program*>(target);
+    };
+#if TINT_BUILD_IR
+    auto get_ir = [&]() {
+        if (std::holds_alternative<const Program*>(target)) {
+            // Convert the AST program to an IR module.
+            auto converted = ir::FromProgram(std::get<const Program*>(target));
+            TINT_ASSERT(Transform, converted);
+            ir_result = converted.Move();
+            target = &ir_result;
+        }
+        TINT_ASSERT(Transform, std::holds_alternative<ir::Module*>(target));
+        return std::get<ir::Module*>(target);
+    };
+#endif  // TINT_BUILD_IR
 
-    TINT_IF_PRINT_PROGRAM(print_program("Input of", this));
+    TINT_IF_PRINT_PROGRAM(print_program("Input of", "transform manager"));
 
     for (const auto& transform : transforms_) {
         if (auto* ast_transform = transform->As<ast::transform::Transform>()) {
-            if (auto result = ast_transform->Apply(program, inputs, outputs)) {
-                output.emplace(std::move(result.value()));
-                program = &output.value();
+            if (auto result = ast_transform->Apply(get_ast(), inputs, outputs)) {
+                ast_result = std::move(result.value());
+                target = &ast_result;
 
-                if (!program->IsValid()) {
-                    TINT_IF_PRINT_PROGRAM(print_program("Invalid output of", transform.get()));
+                if (!ast_result.IsValid()) {
+                    TINT_IF_PRINT_PROGRAM(
+                        print_program("Invalid output of", transform->TypeInfo().name));
                     break;
                 }
 
-                TINT_IF_PRINT_PROGRAM(print_program("Output of", transform.get()));
+                TINT_IF_PRINT_PROGRAM(print_program("Output of", transform->TypeInfo().name));
             } else {
                 TINT_IF_PRINT_PROGRAM(std::cout << "Skipped " << transform->TypeInfo().name
                                                 << std::endl);
             }
+#if TINT_BUILD_IR
+        } else if (auto* ir_transform = transform->As<ir::transform::Transform>()) {
+            ir_transform->Run(get_ir(), inputs, outputs);
+            TINT_IF_PRINT_PROGRAM(print_program("Output of", transform->TypeInfo().name));
+#endif  // TINT_BUILD_IR
         } else {
-            ProgramBuilder b;
-            TINT_ICE(Transform, b.Diagnostics()) << "unhandled transform type";
-            return Program(std::move(b));
+            TINT_ASSERT(Transform, false && "unhandled transform type");
         }
     }
 
-    TINT_IF_PRINT_PROGRAM(print_program("Final output of", this));
+    TINT_IF_PRINT_PROGRAM(print_program("Final output of", "transform manager"));
 
-    if (!output) {
-        ProgramBuilder b;
-        CloneContext ctx{&b, program, /* auto_clone_symbols */ true};
-        ctx.Clone();
-        output = Program(std::move(b));
+    if constexpr (std::is_same<OUTPUT, Program>()) {
+        auto* result = get_ast();
+        if (result == in) {
+            // AST transform pipelines are expected to return a clone of the program, so make sure
+            // the input is cloned at least once even if nothing changed.
+            ProgramBuilder b;
+            CloneContext ctx{&b, result, /* auto_clone_symbols */ true};
+            ctx.Clone();
+            ast_result = Program(std::move(b));
+        }
+        return ast_result;
+#if TINT_BUILD_IR
+    } else if constexpr (std::is_same<OUTPUT, ir::Module*>()) {
+        auto* result = get_ir();
+        if (result == &ir_result) {
+            // IR transform pipelines are expected to mutate the module in place, so move the local
+            // temporary result to the original input.
+            *in = std::move(ir_result);
+        }
+        return in;
+#endif  // TINT_BUILD_IR
     }
-
-    return std::move(output.value());
 }
 
+Program Manager::Run(const Program* program,
+                     const transform::DataMap& inputs,
+                     transform::DataMap& outputs) const {
+    return RunTransforms<Program>(program, inputs, outputs);
+}
+
+#if TINT_BUILD_IR
+void Manager::Run(ir::Module* module, const DataMap& inputs, DataMap& outputs) const {
+    auto* output = RunTransforms<ir::Module*>(module, inputs, outputs);
+    TINT_ASSERT(Transform, output == module);
+}
+#endif  // TINT_BUILD_IR
+
 }  // namespace tint::transform
diff --git a/src/tint/transform/manager.h b/src/tint/transform/manager.h
index 2df5785..289a16d 100644
--- a/src/tint/transform/manager.h
+++ b/src/tint/transform/manager.h
@@ -19,8 +19,16 @@
 #include <utility>
 #include <vector>
 
+#include "src/tint/program.h"
 #include "src/tint/transform/transform.h"
 
+#if TINT_BUILD_IR
+// Forward declarations
+namespace tint::ir {
+class Module;
+}  // namespace tint::ir
+#endif  // TINT_BUILD_IR
+
 namespace tint::transform {
 
 /// A collection of Transforms that act as a single Transform.
@@ -54,8 +62,19 @@
     /// @returns the transformed program
     Program Run(const Program* program, const DataMap& inputs, DataMap& outputs) const;
 
+#if TINT_BUILD_IR
+    /// Runs the transforms on @p module
+    /// @param module the module to transform
+    /// @param inputs optional extra transform-specific input data
+    /// @param outputs optional extra transform-specific output data
+    void Run(ir::Module* module, const DataMap& inputs, DataMap& outputs) const;
+#endif  // TINT_BUILD_IR
+
   private:
     std::vector<std::unique_ptr<Transform>> transforms_;
+
+    template <typename OUTPUT, typename INPUT>
+    OUTPUT RunTransforms(INPUT in, const DataMap& inputs, DataMap& outputs) const;
 };
 
 }  // namespace tint::transform
diff --git a/src/tint/transform/manager_test.cc b/src/tint/transform/manager_test.cc
new file mode 100644
index 0000000..3b0ee25
--- /dev/null
+++ b/src/tint/transform/manager_test.cc
@@ -0,0 +1,174 @@
+// Copyright 2023 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/tint/transform/manager.h"
+
+#include <string>
+
+#include "gtest/gtest.h"
+#include "src/tint/ast/transform/transform.h"
+#include "src/tint/program_builder.h"
+
+#if TINT_BUILD_IR
+#include "src/tint/ir/builder.h"              // nogncheck
+#include "src/tint/ir/transform/transform.h"  // nogncheck
+#endif                                        // TINT_BUILD_IR
+
+namespace tint::transform {
+namespace {
+
+using TransformManagerTest = testing::Test;
+
+class AST_NoOp final : public ast::transform::Transform {
+    ApplyResult Apply(const Program*, const DataMap&, DataMap&) const override {
+        return SkipTransform;
+    }
+};
+
+class AST_AddFunction final : public ast::transform::Transform {
+    ApplyResult Apply(const Program* src, const DataMap&, DataMap&) const override {
+        ProgramBuilder b;
+        CloneContext ctx{&b, src};
+        b.Func(b.Sym("ast_func"), {}, b.ty.void_(), {});
+        ctx.Clone();
+        return Program(std::move(b));
+    }
+};
+
+#if TINT_BUILD_IR
+class IR_AddFunction final : public ir::transform::Transform {
+    void Run(ir::Module* mod, const DataMap&, DataMap&) const override {
+        ir::Builder builder(*mod);
+        auto* func =
+            builder.CreateFunction(mod->symbols.New("ir_func"), mod->types.Get<type::Void>());
+        builder.Branch(func->start_target, func->end_target);
+        mod->functions.Push(func);
+    }
+};
+#endif  // TINT_BUILD_IR
+
+Program MakeAST() {
+    ProgramBuilder b;
+    b.Func(b.Sym("main"), {}, b.ty.void_(), {});
+    return Program(std::move(b));
+}
+
+#if TINT_BUILD_IR
+ir::Module MakeIR() {
+    ir::Module mod;
+    ir::Builder builder(mod);
+    auto* func =
+        builder.CreateFunction(builder.ir.symbols.New("main"), builder.ir.types.Get<type::Void>());
+    builder.Branch(func->start_target, func->end_target);
+    builder.ir.functions.Push(func);
+    return mod;
+}
+#endif  // TINT_BUILD_IR
+
+// Test that an AST program is always cloned, even if all transforms are skipped.
+TEST_F(TransformManagerTest, AST_AlwaysClone) {
+    Program ast = MakeAST();
+
+    transform::Manager manager;
+    transform::DataMap outputs;
+    manager.Add<AST_NoOp>();
+
+    auto result = manager.Run(&ast, {}, outputs);
+    EXPECT_TRUE(result.IsValid()) << result.Diagnostics();
+    EXPECT_NE(result.ID(), ast.ID());
+    ASSERT_EQ(result.AST().Functions().Length(), 1u);
+    EXPECT_EQ(result.AST().Functions()[0]->name->symbol.Name(), "main");
+}
+
+#if TINT_BUILD_IR
+
+// Test that an IR module is mutated in place.
+TEST_F(TransformManagerTest, IR_MutateInPlace) {
+    ir::Module ir = MakeIR();
+
+    transform::Manager manager;
+    transform::DataMap outputs;
+    manager.Add<IR_AddFunction>();
+
+    manager.Run(&ir, {}, outputs);
+    ASSERT_EQ(ir.functions.Length(), 2u);
+    EXPECT_EQ(ir.functions[0]->name.Name(), "main");
+    EXPECT_EQ(ir.functions[1]->name.Name(), "ir_func");
+}
+
+TEST_F(TransformManagerTest, AST_MixedTransforms_AST_Before_IR) {
+    Program ast = MakeAST();
+
+    transform::Manager manager;
+    transform::DataMap outputs;
+    manager.Add<AST_AddFunction>();
+    manager.Add<IR_AddFunction>();
+
+    auto result = manager.Run(&ast, {}, outputs);
+    ASSERT_TRUE(result.IsValid()) << result.Diagnostics();
+    ASSERT_EQ(result.AST().Functions().Length(), 3u);
+    EXPECT_EQ(result.AST().Functions()[0]->name->symbol.Name(), "ast_func");
+    EXPECT_EQ(result.AST().Functions()[1]->name->symbol.Name(), "main");
+    EXPECT_EQ(result.AST().Functions()[2]->name->symbol.Name(), "ir_func");
+}
+
+TEST_F(TransformManagerTest, AST_MixedTransforms_IR_Before_AST) {
+    Program ast = MakeAST();
+
+    transform::Manager manager;
+    transform::DataMap outputs;
+    manager.Add<IR_AddFunction>();
+    manager.Add<AST_AddFunction>();
+
+    auto result = manager.Run(&ast, {}, outputs);
+    ASSERT_TRUE(result.IsValid()) << result.Diagnostics();
+    ASSERT_EQ(result.AST().Functions().Length(), 3u);
+    EXPECT_EQ(result.AST().Functions()[0]->name->symbol.Name(), "ast_func");
+    EXPECT_EQ(result.AST().Functions()[1]->name->symbol.Name(), "main");
+    EXPECT_EQ(result.AST().Functions()[2]->name->symbol.Name(), "ir_func");
+}
+
+TEST_F(TransformManagerTest, IR_MixedTransforms_AST_Before_IR) {
+    ir::Module ir = MakeIR();
+
+    transform::Manager manager;
+    transform::DataMap outputs;
+    manager.Add<AST_AddFunction>();
+    manager.Add<IR_AddFunction>();
+
+    manager.Run(&ir, {}, outputs);
+    ASSERT_EQ(ir.functions.Length(), 3u);
+    EXPECT_EQ(ir.functions[0]->name.Name(), "ast_func");
+    EXPECT_EQ(ir.functions[1]->name.Name(), "main");
+    EXPECT_EQ(ir.functions[2]->name.Name(), "ir_func");
+}
+
+TEST_F(TransformManagerTest, IR_MixedTransforms_IR_Before_AST) {
+    ir::Module ir = MakeIR();
+
+    transform::Manager manager;
+    transform::DataMap outputs;
+    manager.Add<IR_AddFunction>();
+    manager.Add<AST_AddFunction>();
+
+    manager.Run(&ir, {}, outputs);
+    ASSERT_EQ(ir.functions.Length(), 3u);
+    EXPECT_EQ(ir.functions[0]->name.Name(), "ast_func");
+    EXPECT_EQ(ir.functions[1]->name.Name(), "main");
+    EXPECT_EQ(ir.functions[2]->name.Name(), "ir_func");
+}
+#endif  // TINT_BUILD_IR
+
+}  // namespace
+}  // namespace tint::transform
diff --git a/src/tint/transform/transform.cc b/src/tint/transform/transform.cc
index 657f9be..5d2117c 100644
--- a/src/tint/transform/transform.cc
+++ b/src/tint/transform/transform.cc
@@ -14,8 +14,6 @@
 
 #include "src/tint/transform/transform.h"
 
-#include "src/tint/program_builder.h"
-
 TINT_INSTANTIATE_TYPEINFO(tint::transform::Transform);
 TINT_INSTANTIATE_TYPEINFO(tint::transform::Data);
 
diff --git a/src/tint/transform/transform.h b/src/tint/transform/transform.h
index 1f8a3f5..d142a96 100644
--- a/src/tint/transform/transform.h
+++ b/src/tint/transform/transform.h
@@ -19,7 +19,6 @@
 #include <unordered_map>
 #include <utility>
 
-#include "src/tint/program.h"
 #include "src/tint/utils/castable.h"
 
 namespace tint::transform {
diff --git a/src/tint/writer/spirv/generator.cc b/src/tint/writer/spirv/generator.cc
index de6178b..6531ec5 100644
--- a/src/tint/writer/spirv/generator.cc
+++ b/src/tint/writer/spirv/generator.cc
@@ -18,9 +18,9 @@
 
 #include "src/tint/writer/spirv/generator_impl.h"
 #if TINT_BUILD_IR
-#include "src/tint/ir/from_program.h"                 // nogncheck
-#include "src/tint/writer/spirv/generator_impl_ir.h"  // nogncheck
-#endif                                                // TINT_BUILD_IR
+#include "src/tint/ir/from_program.h"                    // nogncheck
+#include "src/tint/writer/spirv/ir/generator_impl_ir.h"  // nogncheck
+#endif                                                   // TINT_BUILD_IR
 
 namespace tint::writer::spirv {
 
@@ -41,14 +41,15 @@
 #if TINT_BUILD_IR
     if (options.use_tint_ir) {
         // Convert the AST program to an IR module.
-        auto ir = ir::FromProgram(program);
-        if (!ir) {
-            result.error = "IR converter: " + ir.Failure();
+        auto converted = ir::FromProgram(program);
+        if (!converted) {
+            result.error = "IR converter: " + converted.Failure();
             return result;
         }
 
         // Generate the SPIR-V code.
-        auto impl = std::make_unique<GeneratorImplIr>(&ir.Get(), zero_initialize_workgroup_memory);
+        auto ir = converted.Move();
+        auto impl = std::make_unique<GeneratorImplIr>(&ir, zero_initialize_workgroup_memory);
         result.success = impl->Generate();
         result.error = impl->Diagnostics().str();
         result.spirv = std::move(impl->Result());
diff --git a/src/tint/writer/spirv/generator_impl_ir.cc b/src/tint/writer/spirv/ir/generator_impl_ir.cc
similarity index 67%
rename from src/tint/writer/spirv/generator_impl_ir.cc
rename to src/tint/writer/spirv/ir/generator_impl_ir.cc
index 5d8c2f0..4ec4d9d 100644
--- a/src/tint/writer/spirv/generator_impl_ir.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.cc
@@ -12,30 +12,71 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "src/tint/writer/spirv/generator_impl_ir.h"
+#include "src/tint/writer/spirv/ir/generator_impl_ir.h"
 
 #include "spirv/unified1/spirv.h"
 #include "src/tint/ir/binary.h"
 #include "src/tint/ir/block.h"
 #include "src/tint/ir/function_terminator.h"
+#include "src/tint/ir/if.h"
 #include "src/tint/ir/module.h"
+#include "src/tint/ir/store.h"
+#include "src/tint/ir/transform/add_empty_entry_point.h"
+#include "src/tint/ir/var.h"
 #include "src/tint/switch.h"
+#include "src/tint/transform/manager.h"
 #include "src/tint/type/bool.h"
 #include "src/tint/type/f16.h"
 #include "src/tint/type/f32.h"
 #include "src/tint/type/i32.h"
+#include "src/tint/type/pointer.h"
 #include "src/tint/type/type.h"
 #include "src/tint/type/u32.h"
 #include "src/tint/type/vector.h"
 #include "src/tint/type/void.h"
+#include "src/tint/writer/spirv/generator.h"
 #include "src/tint/writer/spirv/module.h"
 
 namespace tint::writer::spirv {
 
-GeneratorImplIr::GeneratorImplIr(const ir::Module* module, bool zero_init_workgroup_mem)
+namespace {
+
+void Sanitize(ir::Module* module) {
+    transform::Manager manager;
+    transform::DataMap data;
+
+    manager.Add<ir::transform::AddEmptyEntryPoint>();
+
+    transform::DataMap outputs;
+    manager.Run(module, data, outputs);
+}
+
+SpvStorageClass StorageClass(builtin::AddressSpace addrspace) {
+    switch (addrspace) {
+        case builtin::AddressSpace::kFunction:
+            return SpvStorageClassFunction;
+        case builtin::AddressSpace::kPrivate:
+            return SpvStorageClassPrivate;
+        case builtin::AddressSpace::kStorage:
+            return SpvStorageClassStorageBuffer;
+        case builtin::AddressSpace::kUniform:
+            return SpvStorageClassUniform;
+        case builtin::AddressSpace::kWorkgroup:
+            return SpvStorageClassWorkgroup;
+        default:
+            return SpvStorageClassMax;
+    }
+}
+
+}  // namespace
+
+GeneratorImplIr::GeneratorImplIr(ir::Module* module, bool zero_init_workgroup_mem)
     : ir_(module), zero_init_workgroup_memory_(zero_init_workgroup_mem) {}
 
 bool GeneratorImplIr::Generate() {
+    // Run the IR transformations to prepare for SPIR-V emission.
+    Sanitize(ir_);
+
     // TODO(crbug.com/tint/1906): Check supported extensions.
 
     module_.PushCapability(SpvCapabilityShader);
@@ -133,6 +174,11 @@
             [&](const type::Vector* vec) {
                 module_.PushType(spv::Op::OpTypeVector, {id, Type(vec->type()), vec->Width()});
             },
+            [&](const type::Pointer* ptr) {
+                module_.PushType(
+                    spv::Op::OpTypePointer,
+                    {id, U32Operand(StorageClass(ptr->AddressSpace())), Type(ptr->StoreType())});
+            },
             [&](Default) {
                 TINT_ICE(Writer, diagnostics_) << "unhandled type: " << ty->FriendlyName();
             });
@@ -158,6 +204,10 @@
         });
 }
 
+uint32_t GeneratorImplIr::Label(const ir::Block* block) {
+    return block_labels_.GetOrCreate(block, [&]() { return module_.NextId(); });
+}
+
 void GeneratorImplIr::EmitFunction(const ir::Function* func) {
     // Make an ID for the function.
     auto id = module_.NextId();
@@ -235,11 +285,22 @@
 }
 
 void GeneratorImplIr::EmitBlock(const ir::Block* block) {
+    // Emit the label.
+    // Skip if this is the function's entry block, as it will be emitted by the function object.
+    if (!current_function_.instructions().empty()) {
+        current_function_.push_inst(spv::Op::OpLabel, {Label(block)});
+    }
+
     // Emit the instructions.
     for (auto* inst : block->instructions) {
         auto result = Switch(
             inst,  //
             [&](const ir::Binary* b) { return EmitBinary(b); },
+            [&](const ir::Store* s) {
+                EmitStore(s);
+                return 0u;
+            },
+            [&](const ir::Var* v) { return EmitVar(v); },
             [&](Default) {
                 TINT_ICE(Writer, diagnostics_)
                     << "unimplemented instruction: " << inst->TypeInfo().name;
@@ -251,6 +312,8 @@
     // Handle the branch at the end of the block.
     Switch(
         block->branch.target,
+        [&](const ir::Block* b) { current_function_.push_inst(spv::Op::OpBranch, {Label(b)}); },
+        [&](const ir::If* i) { EmitIf(i); },
         [&](const ir::FunctionTerminator*) {
             // TODO(jrprice): Handle the return value, which will be a branch argument.
             if (!block->branch.args.IsEmpty()) {
@@ -258,7 +321,52 @@
             }
             current_function_.push_inst(spv::Op::OpReturn, {});
         },
-        [&](Default) { TINT_ICE(Writer, diagnostics_) << "unimplemented branch target"; });
+        [&](Default) {
+            if (!block->branch.target) {
+                // A block may not have an outward branch (e.g. an unreachable merge block).
+                current_function_.push_inst(spv::Op::OpUnreachable, {});
+            } else {
+                TINT_ICE(Writer, diagnostics_)
+                    << "unimplemented branch target: " << block->branch.target->TypeInfo().name;
+            }
+        });
+}
+
+void GeneratorImplIr::EmitIf(const ir::If* i) {
+    auto* merge_block = i->merge.target->As<ir::Block>();
+    auto* true_block = i->true_.target->As<ir::Block>();
+    auto* false_block = i->false_.target->As<ir::Block>();
+
+    // Generate labels for the blocks. We emit the true or false block if it:
+    // 1. contains instructions, or
+    // 2. branches somewhere other then the merge target.
+    // Otherwise we skip them and branch straight to the merge block.
+    uint32_t merge_label = Label(merge_block);
+    uint32_t true_label = merge_label;
+    uint32_t false_label = merge_label;
+    if (!true_block->instructions.IsEmpty() || true_block->branch.target != merge_block) {
+        true_label = Label(true_block);
+    }
+    if (!false_block->instructions.IsEmpty() || false_block->branch.target != merge_block) {
+        false_label = Label(false_block);
+    }
+
+    // Emit the OpSelectionMerge and OpBranchConditional instructions.
+    current_function_.push_inst(spv::Op::OpSelectionMerge,
+                                {merge_label, U32Operand(SpvSelectionControlMaskNone)});
+    current_function_.push_inst(spv::Op::OpBranchConditional,
+                                {Value(i->condition), true_label, false_label});
+
+    // Emit the `true` and `false` blocks, if they're not being skipped.
+    if (true_label != merge_label) {
+        EmitBlock(true_block);
+    }
+    if (false_label != merge_label) {
+        EmitBlock(false_block);
+    }
+
+    // Emit the merge block.
+    EmitBlock(merge_block);
 }
 
 uint32_t GeneratorImplIr::EmitBinary(const ir::Binary* binary) {
@@ -288,4 +396,34 @@
     return id;
 }
 
+void GeneratorImplIr::EmitStore(const ir::Store* store) {
+    current_function_.push_inst(spv::Op::OpStore, {Value(store->to), Value(store->from)});
+}
+
+uint32_t GeneratorImplIr::EmitVar(const ir::Var* var) {
+    auto id = module_.NextId();
+    auto* ptr = var->Type()->As<type::Pointer>();
+    TINT_ASSERT(Writer, ptr);
+    auto ty = Type(ptr);
+
+    if (ptr->AddressSpace() == builtin::AddressSpace::kFunction) {
+        TINT_ASSERT(Writer, current_function_);
+        current_function_.push_var({ty, id, U32Operand(SpvStorageClassFunction)});
+        if (var->initializer) {
+            current_function_.push_inst(spv::Op::OpStore, {id, Value(var->initializer)});
+        }
+    } else {
+        TINT_ICE(Writer, diagnostics_)
+            << "unimplemented variable address space " << ptr->AddressSpace();
+        return 0u;
+    }
+
+    // Set the name if present.
+    if (auto name = ir_->NameOf(var)) {
+        module_.PushDebug(spv::Op::OpName, {id, Operand(name.Name())});
+    }
+
+    return id;
+}
+
 }  // namespace tint::writer::spirv
diff --git a/src/tint/writer/spirv/generator_impl_ir.h b/src/tint/writer/spirv/ir/generator_impl_ir.h
similarity index 85%
rename from src/tint/writer/spirv/generator_impl_ir.h
rename to src/tint/writer/spirv/ir/generator_impl_ir.h
index e885edb..66ffe48 100644
--- a/src/tint/writer/spirv/generator_impl_ir.h
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.h
@@ -12,8 +12,8 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#ifndef SRC_TINT_WRITER_SPIRV_GENERATOR_IMPL_IR_H_
-#define SRC_TINT_WRITER_SPIRV_GENERATOR_IMPL_IR_H_
+#ifndef SRC_TINT_WRITER_SPIRV_IR_GENERATOR_IMPL_IR_H_
+#define SRC_TINT_WRITER_SPIRV_IR_GENERATOR_IMPL_IR_H_
 
 #include <vector>
 
@@ -30,9 +30,12 @@
 namespace tint::ir {
 class Binary;
 class Block;
+class If;
 class Function;
 class Module;
+class Store;
 class Value;
+class Var;
 }  // namespace tint::ir
 namespace tint::type {
 class Type;
@@ -47,7 +50,7 @@
     /// @param module the Tint IR module to generate
     /// @param zero_init_workgroup_memory `true` to initialize all the variables in the Workgroup
     ///                                   storage class with OpConstantNull
-    GeneratorImplIr(const ir::Module* module, bool zero_init_workgroup_memory);
+    GeneratorImplIr(ir::Module* module, bool zero_init_workgroup_memory);
 
     /// @returns true on successful generation; false otherwise
     bool Generate();
@@ -76,6 +79,11 @@
     /// @returns the result ID of the value
     uint32_t Value(const ir::Value* value);
 
+    /// Get the ID of the label for `block`.
+    /// @param block the block to get the label ID for
+    /// @returns the ID of the block's label
+    uint32_t Label(const ir::Block* block);
+
     /// Emit a function.
     /// @param func the function to emit
     void EmitFunction(const ir::Function* func);
@@ -89,18 +97,31 @@
     /// @param block the block to emit
     void EmitBlock(const ir::Block* block);
 
+    /// Emit an `if` flow node.
+    /// @param i the if node to emit
+    void EmitIf(const ir::If* i);
+
     /// Emit a binary instruction.
     /// @param binary the binary instruction to emit
     /// @returns the result ID of the instruction
     uint32_t EmitBinary(const ir::Binary* binary);
 
+    /// Emit a store instruction.
+    /// @param store the store instruction to emit
+    void EmitStore(const ir::Store* store);
+
+    /// Emit a var instruction.
+    /// @param var the var instruction to emit
+    /// @returns the result ID of the instruction
+    uint32_t EmitVar(const ir::Var* var);
+
   private:
     /// Get the result ID of the constant `constant`, emitting its instruction if necessary.
     /// @param constant the constant to get the ID for
     /// @returns the result ID of the constant
     uint32_t Constant(const constant::Value* constant);
 
-    const ir::Module* ir_;
+    ir::Module* ir_;
     spirv::Module module_;
     BinaryWriter writer_;
     diag::List diagnostics_;
@@ -161,6 +182,9 @@
     /// The map of instructions to their result IDs.
     utils::Hashmap<const ir::Instruction*, uint32_t, 8> instructions_;
 
+    /// The map of blocks to the IDs of their label instructions.
+    utils::Hashmap<const ir::Block*, uint32_t, 8> block_labels_;
+
     /// The current function that is being emitted.
     Function current_function_;
 
@@ -169,4 +193,4 @@
 
 }  // namespace tint::writer::spirv
 
-#endif  // SRC_TINT_WRITER_SPIRV_GENERATOR_IMPL_IR_H_
+#endif  // SRC_TINT_WRITER_SPIRV_IR_GENERATOR_IMPL_IR_H_
diff --git a/src/tint/writer/spirv/generator_impl_binary_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
similarity index 98%
rename from src/tint/writer/spirv/generator_impl_binary_test.cc
rename to src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
index 4bf6b81..30dc059 100644
--- a/src/tint/writer/spirv/generator_impl_binary_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
@@ -12,7 +12,7 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "src/tint/writer/spirv/test_helper_ir.h"
+#include "src/tint/writer/spirv/ir/test_helper_ir.h"
 
 using namespace tint::number_suffixes;  // NOLINT
 
diff --git a/src/tint/writer/spirv/generator_impl_constant_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc
similarity index 98%
rename from src/tint/writer/spirv/generator_impl_constant_test.cc
rename to src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc
index 1e27e89..95fce03 100644
--- a/src/tint/writer/spirv/generator_impl_constant_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc
@@ -12,7 +12,7 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "src/tint/writer/spirv/test_helper_ir.h"
+#include "src/tint/writer/spirv/ir/test_helper_ir.h"
 
 namespace tint::writer::spirv {
 namespace {
diff --git a/src/tint/writer/spirv/generator_impl_function_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc
similarity index 98%
rename from src/tint/writer/spirv/generator_impl_function_test.cc
rename to src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc
index b8faebb..77b4a62 100644
--- a/src/tint/writer/spirv/generator_impl_function_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc
@@ -12,7 +12,7 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "src/tint/writer/spirv/test_helper_ir.h"
+#include "src/tint/writer/spirv/ir/test_helper_ir.h"
 
 namespace tint::writer::spirv {
 namespace {
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
new file mode 100644
index 0000000..7b41184
--- /dev/null
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
@@ -0,0 +1,149 @@
+// Copyright 2023 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/tint/writer/spirv/ir/test_helper_ir.h"
+
+using namespace tint::number_suffixes;  // NOLINT
+
+namespace tint::writer::spirv {
+namespace {
+
+TEST_F(SpvGeneratorImplTest, If_TrueEmpty_FalseEmpty) {
+    auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
+
+    auto* i = b.CreateIf(b.Constant(true));
+    b.Branch(i->true_.target->As<ir::Block>(), i->merge.target);
+    b.Branch(i->false_.target->As<ir::Block>(), i->merge.target);
+    b.Branch(i->merge.target->As<ir::Block>(), func->end_target);
+
+    b.Branch(func->start_target, i);
+
+    generator_.EmitFunction(func);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%7 = OpTypeBool
+%6 = OpConstantTrue %7
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpSelectionMerge %5 None
+OpBranchConditional %6 %5 %5
+%5 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, If_FalseEmpty) {
+    auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
+
+    auto* i = b.CreateIf(b.Constant(true));
+    b.Branch(i->false_.target->As<ir::Block>(), i->merge.target);
+    b.Branch(i->merge.target->As<ir::Block>(), func->end_target);
+
+    auto* true_block = i->true_.target->As<ir::Block>();
+    true_block->instructions.Push(
+        b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(1_i)));
+    b.Branch(true_block, i->merge.target);
+
+    b.Branch(func->start_target, i);
+
+    generator_.EmitFunction(func);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%8 = OpTypeBool
+%7 = OpConstantTrue %8
+%10 = OpTypeInt 32 1
+%11 = OpConstant %10 1
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpSelectionMerge %5 None
+OpBranchConditional %7 %6 %5
+%6 = OpLabel
+%9 = OpIAdd %10 %11 %11
+OpBranch %5
+%5 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, If_TrueEmpty) {
+    auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
+
+    auto* i = b.CreateIf(b.Constant(true));
+    b.Branch(i->true_.target->As<ir::Block>(), i->merge.target);
+    b.Branch(i->merge.target->As<ir::Block>(), func->end_target);
+
+    auto* false_block = i->false_.target->As<ir::Block>();
+    false_block->instructions.Push(
+        b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(1_i)));
+    b.Branch(false_block, i->merge.target);
+
+    b.Branch(func->start_target, i);
+
+    generator_.EmitFunction(func);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%8 = OpTypeBool
+%7 = OpConstantTrue %8
+%10 = OpTypeInt 32 1
+%11 = OpConstant %10 1
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpSelectionMerge %5 None
+OpBranchConditional %7 %5 %6
+%6 = OpLabel
+%9 = OpIAdd %10 %11 %11
+OpBranch %5
+%5 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, If_BothBranchesReturn) {
+    auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
+
+    auto* i = b.CreateIf(b.Constant(true));
+    b.Branch(i->true_.target->As<ir::Block>(), func->end_target);
+    b.Branch(i->false_.target->As<ir::Block>(), func->end_target);
+    i->merge.target->As<ir::Block>()->branch.target = nullptr;
+
+    b.Branch(func->start_target, i);
+
+    generator_.EmitFunction(func);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%9 = OpTypeBool
+%8 = OpConstantTrue %9
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpSelectionMerge %5 None
+OpBranchConditional %8 %6 %7
+%6 = OpLabel
+OpReturn
+%7 = OpLabel
+OpReturn
+%5 = OpLabel
+OpUnreachable
+OpFunctionEnd
+)");
+}
+
+}  // namespace
+}  // namespace tint::writer::spirv
diff --git a/src/tint/writer/spirv/generator_impl_ir_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_test.cc
similarity index 85%
rename from src/tint/writer/spirv/generator_impl_ir_test.cc
rename to src/tint/writer/spirv/ir/generator_impl_ir_test.cc
index a202eea..ffafa60 100644
--- a/src/tint/writer/spirv/generator_impl_ir_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_test.cc
@@ -12,7 +12,9 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "src/tint/writer/spirv/test_helper_ir.h"
+#include "src/tint/writer/spirv/ir/test_helper_ir.h"
+
+#include "gmock/gmock.h"
 
 namespace tint::writer::spirv {
 namespace {
@@ -20,9 +22,9 @@
 TEST_F(SpvGeneratorImplTest, ModuleHeader) {
     ASSERT_TRUE(generator_.Generate()) << generator_.Diagnostics().str();
     auto got = Disassemble(generator_.Result());
-    EXPECT_EQ(got, R"(OpCapability Shader
+    EXPECT_THAT(got, testing::StartsWith(R"(OpCapability Shader
 OpMemoryModel Logical GLSL450
-)");
+)"));
 }
 
 }  // namespace
diff --git a/src/tint/writer/spirv/generator_impl_type_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_type_test.cc
similarity index 98%
rename from src/tint/writer/spirv/generator_impl_type_test.cc
rename to src/tint/writer/spirv/ir/generator_impl_ir_type_test.cc
index c86d363..38bf1f0 100644
--- a/src/tint/writer/spirv/generator_impl_type_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_type_test.cc
@@ -19,7 +19,7 @@
 #include "src/tint/type/type.h"
 #include "src/tint/type/u32.h"
 #include "src/tint/type/void.h"
-#include "src/tint/writer/spirv/test_helper_ir.h"
+#include "src/tint/writer/spirv/ir/test_helper_ir.h"
 
 namespace tint::writer::spirv {
 namespace {
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
new file mode 100644
index 0000000..ff94862
--- /dev/null
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
@@ -0,0 +1,168 @@
+// Copyright 2023 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/tint/type/pointer.h"
+#include "src/tint/writer/spirv/ir/test_helper_ir.h"
+
+using namespace tint::number_suffixes;  // NOLINT
+
+namespace tint::writer::spirv {
+namespace {
+
+TEST_F(SpvGeneratorImplTest, FunctionVar_NoInit) {
+    auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
+    b.Branch(func->start_target, func->end_target);
+
+    auto* ty = mod.types.Get<type::Pointer>(
+        mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
+    auto* v = b.Declare(ty);
+    func->start_target->instructions.Push(v);
+
+    generator_.EmitFunction(func);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%7 = OpTypeInt 32 1
+%6 = OpTypePointer Function %7
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+%5 = OpVariable %6 Function
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, FunctionVar_WithInit) {
+    auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
+    b.Branch(func->start_target, func->end_target);
+
+    auto* ty = mod.types.Get<type::Pointer>(
+        mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
+    auto* v = b.Declare(ty);
+    func->start_target->instructions.Push(v);
+    v->initializer = b.Constant(42_i);
+
+    generator_.EmitFunction(func);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%7 = OpTypeInt 32 1
+%6 = OpTypePointer Function %7
+%8 = OpConstant %7 42
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+%5 = OpVariable %6 Function
+OpStore %5 %8
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, FunctionVar_Name) {
+    auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
+    b.Branch(func->start_target, func->end_target);
+
+    auto* ty = mod.types.Get<type::Pointer>(
+        mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
+    auto* v = b.Declare(ty);
+    func->start_target->instructions.Push(v);
+    mod.SetName(v, "myvar");
+
+    generator_.EmitFunction(func);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+OpName %5 "myvar"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%7 = OpTypeInt 32 1
+%6 = OpTypePointer Function %7
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+%5 = OpVariable %6 Function
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, FunctionVar_DeclInsideBlock) {
+    auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
+    b.Branch(func->start_target, func->end_target);
+
+    auto* ty = mod.types.Get<type::Pointer>(
+        mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
+    auto* v = b.Declare(ty);
+    v->initializer = b.Constant(42_i);
+
+    auto* i = b.CreateIf(b.Constant(true));
+    b.Branch(i->false_.target->As<ir::Block>(), func->end_target);
+    b.Branch(i->merge.target->As<ir::Block>(), func->end_target);
+
+    auto* true_block = i->true_.target->As<ir::Block>();
+    true_block->instructions.Push(v);
+    b.Branch(true_block, i->merge.target);
+
+    b.Branch(func->start_target, i);
+
+    generator_.EmitFunction(func);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%9 = OpTypeBool
+%8 = OpConstantTrue %9
+%12 = OpTypeInt 32 1
+%11 = OpTypePointer Function %12
+%13 = OpConstant %12 42
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+%10 = OpVariable %11 Function
+OpSelectionMerge %5 None
+OpBranchConditional %8 %6 %7
+%6 = OpLabel
+OpStore %10 %13
+OpBranch %5
+%7 = OpLabel
+OpReturn
+%5 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, FunctionVar_Store) {
+    auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
+    b.Branch(func->start_target, func->end_target);
+
+    auto* ty = mod.types.Get<type::Pointer>(
+        mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
+    auto* v = b.Declare(ty);
+    func->start_target->instructions.Push(v);
+    func->start_target->instructions.Push(b.Store(v, b.Constant(42_i)));
+
+    generator_.EmitFunction(func);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%7 = OpTypeInt 32 1
+%6 = OpTypePointer Function %7
+%8 = OpConstant %7 42
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+%5 = OpVariable %6 Function
+OpStore %5 %8
+OpReturn
+OpFunctionEnd
+)");
+}
+
+}  // namespace
+}  // namespace tint::writer::spirv
diff --git a/src/tint/writer/spirv/test_helper_ir.h b/src/tint/writer/spirv/ir/test_helper_ir.h
similarity index 87%
rename from src/tint/writer/spirv/test_helper_ir.h
rename to src/tint/writer/spirv/ir/test_helper_ir.h
index 3574645..9509b42 100644
--- a/src/tint/writer/spirv/test_helper_ir.h
+++ b/src/tint/writer/spirv/ir/test_helper_ir.h
@@ -12,14 +12,14 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#ifndef SRC_TINT_WRITER_SPIRV_TEST_HELPER_IR_H_
-#define SRC_TINT_WRITER_SPIRV_TEST_HELPER_IR_H_
+#ifndef SRC_TINT_WRITER_SPIRV_IR_TEST_HELPER_IR_H_
+#define SRC_TINT_WRITER_SPIRV_IR_TEST_HELPER_IR_H_
 
 #include <string>
 
 #include "gtest/gtest.h"
 #include "src/tint/ir/builder.h"
-#include "src/tint/writer/spirv/generator_impl_ir.h"
+#include "src/tint/writer/spirv/ir/generator_impl_ir.h"
 #include "src/tint/writer/spirv/spv_dump.h"
 
 namespace tint::writer::spirv {
@@ -50,4 +50,4 @@
 
 }  // namespace tint::writer::spirv
 
-#endif  // SRC_TINT_WRITER_SPIRV_TEST_HELPER_IR_H_
+#endif  // SRC_TINT_WRITER_SPIRV_IR_TEST_HELPER_IR_H_