diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 47efd52..64fde4c 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1226,6 +1226,8 @@
     "ir/flow_node.h",
     "ir/function.cc",
     "ir/function.h",
+    "ir/function_param.cc",
+    "ir/function_param.h",
     "ir/function_terminator.cc",
     "ir/function_terminator.h",
     "ir/if.cc",
@@ -2263,6 +2265,7 @@
       "ir/store_test.cc",
       "ir/test_helper.h",
       "ir/to_program_roundtrip_test.cc",
+      "ir/transform/add_empty_entry_point_test.cc",
       "ir/unary_test.cc",
     ]
 
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index d0c7487..8ec573a 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -742,6 +742,8 @@
     ir/flow_node.h
     ir/function.cc
     ir/function.h
+    ir/function_param.cc
+    ir/function_param.h
     ir/function_terminator.cc
     ir/function_terminator.h
     ir/if.cc
@@ -1474,6 +1476,7 @@
       ir/module_test.cc
       ir/store_test.cc
       ir/test_helper.h
+      ir/transform/add_empty_entry_point_test.cc
       ir/unary_test.cc
     )
   endif()
diff --git a/src/tint/ir/binary.cc b/src/tint/ir/binary.cc
index af28306..2b179ac 100644
--- a/src/tint/ir/binary.cc
+++ b/src/tint/ir/binary.cc
@@ -19,8 +19,8 @@
 
 namespace tint::ir {
 
-Binary::Binary(Kind k, const type::Type* res_ty, Value* lhs, Value* rhs)
-    : kind(k), result_type(res_ty), lhs_(lhs), rhs_(rhs) {
+Binary::Binary(enum Kind kind, const type::Type* res_ty, Value* lhs, Value* rhs)
+    : kind_(kind), result_type_(res_ty), lhs_(lhs), rhs_(rhs) {
     TINT_ASSERT(IR, lhs);
     TINT_ASSERT(IR, rhs);
     lhs_->AddUsage(this);
diff --git a/src/tint/ir/binary.h b/src/tint/ir/binary.h
index a6dc722..7ea6c08 100644
--- a/src/tint/ir/binary.h
+++ b/src/tint/ir/binary.h
@@ -51,7 +51,7 @@
     /// @param type the result type
     /// @param lhs the lhs of the instruction
     /// @param rhs the rhs of the instruction
-    Binary(Kind kind, const type::Type* type, Value* lhs, Value* rhs);
+    Binary(enum Kind kind, const type::Type* type, Value* lhs, Value* rhs);
     Binary(const Binary& inst) = delete;
     Binary(Binary&& inst) = delete;
     ~Binary() override;
@@ -59,8 +59,11 @@
     Binary& operator=(const Binary& inst) = delete;
     Binary& operator=(Binary&& inst) = delete;
 
+    /// @returns the kind of the binary instruction
+    enum Kind Kind() const { return kind_; }
+
     /// @returns the type of the value
-    const type::Type* Type() const override { return result_type; }
+    const type::Type* Type() const override { return result_type_; }
 
     /// @returns the left-hand-side value for the instruction
     const Value* LHS() const { return lhs_; }
@@ -68,15 +71,11 @@
     /// @returns the right-hand-side value for the instruction
     const Value* RHS() const { return rhs_; }
 
-    /// the kind of binary instruction
-    Kind kind = Kind::kAdd;
-
-    /// the result type of the instruction
-    const type::Type* result_type = nullptr;
-
   private:
-    Value* lhs_ = nullptr;
-    Value* rhs_ = nullptr;
+    enum Kind kind_;
+    const type::Type* result_type_;
+    Value* lhs_;
+    Value* rhs_;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/binary_test.cc b/src/tint/ir/binary_test.cc
index 2968d28..dc35978 100644
--- a/src/tint/ir/binary_test.cc
+++ b/src/tint/ir/binary_test.cc
@@ -30,17 +30,16 @@
     const auto* inst = b.And(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kAnd);
-    ASSERT_NE(inst->result_type, nullptr);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kAnd);
     ASSERT_NE(inst->Type(), nullptr);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -52,15 +51,15 @@
     const auto* inst = b.Or(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kOr);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kOr);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -72,15 +71,15 @@
     const auto* inst = b.Xor(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kXor);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kXor);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -92,15 +91,15 @@
     const auto* inst = b.Equal(b.ir.types.Get<type::Bool>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kEqual);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kEqual);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -112,15 +111,15 @@
     const auto* inst = b.NotEqual(b.ir.types.Get<type::Bool>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kNotEqual);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kNotEqual);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -132,15 +131,15 @@
     const auto* inst = b.LessThan(b.ir.types.Get<type::Bool>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kLessThan);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kLessThan);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -153,15 +152,15 @@
         b.GreaterThan(b.ir.types.Get<type::Bool>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kGreaterThan);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kGreaterThan);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -174,15 +173,15 @@
         b.LessThanEqual(b.ir.types.Get<type::Bool>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kLessThanEqual);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kLessThanEqual);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -195,15 +194,15 @@
         b.GreaterThanEqual(b.ir.types.Get<type::Bool>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kGreaterThanEqual);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kGreaterThanEqual);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -214,15 +213,15 @@
     const auto* inst = b.Not(b.ir.types.Get<type::Bool>(), b.Constant(true));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kEqual);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kEqual);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<bool>>());
     EXPECT_TRUE(lhs->As<constant::Scalar<bool>>()->ValueAs<bool>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<bool>>());
     EXPECT_FALSE(rhs->As<constant::Scalar<bool>>()->ValueAs<bool>());
 }
@@ -234,15 +233,15 @@
     const auto* inst = b.ShiftLeft(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kShiftLeft);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kShiftLeft);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -254,15 +253,15 @@
     const auto* inst = b.ShiftRight(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kShiftRight);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kShiftRight);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -274,15 +273,15 @@
     const auto* inst = b.Add(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kAdd);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kAdd);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -294,15 +293,15 @@
     const auto* inst = b.Subtract(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kSubtract);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kSubtract);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -314,15 +313,15 @@
     const auto* inst = b.Multiply(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kMultiply);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kMultiply);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -334,15 +333,15 @@
     const auto* inst = b.Divide(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kDivide);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kDivide);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -354,15 +353,15 @@
     const auto* inst = b.Modulo(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
 
     ASSERT_TRUE(inst->Is<Binary>());
-    EXPECT_EQ(inst->kind, Binary::Kind::kModulo);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kModulo);
 
     ASSERT_TRUE(inst->LHS()->Is<Constant>());
-    auto lhs = inst->LHS()->As<Constant>()->value;
+    auto lhs = inst->LHS()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
     ASSERT_TRUE(inst->RHS()->Is<Constant>());
-    auto rhs = inst->RHS()->As<Constant>()->value;
+    auto rhs = inst->RHS()->As<Constant>()->Value();
     ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -372,7 +371,7 @@
     Builder b{mod};
     const auto* inst = b.And(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
 
-    EXPECT_EQ(inst->kind, Binary::Kind::kAnd);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kAnd);
 
     ASSERT_NE(inst->LHS(), nullptr);
     ASSERT_EQ(inst->LHS()->Usage().Length(), 1u);
@@ -389,7 +388,7 @@
     auto val = b.Constant(4_i);
     const auto* inst = b.And(b.ir.types.Get<type::I32>(), val, val);
 
-    EXPECT_EQ(inst->kind, Binary::Kind::kAnd);
+    EXPECT_EQ(inst->Kind(), Binary::Kind::kAnd);
     ASSERT_EQ(inst->LHS(), inst->RHS());
 
     ASSERT_NE(inst->LHS(), nullptr);
diff --git a/src/tint/ir/bitcast_test.cc b/src/tint/ir/bitcast_test.cc
index e78b734..6eda562 100644
--- a/src/tint/ir/bitcast_test.cc
+++ b/src/tint/ir/bitcast_test.cc
@@ -32,9 +32,10 @@
     ASSERT_TRUE(inst->Is<ir::Bitcast>());
     ASSERT_NE(inst->Type(), nullptr);
 
-    ASSERT_EQ(inst->args.Length(), 1u);
-    ASSERT_TRUE(inst->args[0]->Is<Constant>());
-    auto val = inst->args[0]->As<Constant>()->value;
+    const auto args = inst->Args();
+    ASSERT_EQ(args.Length(), 1u);
+    ASSERT_TRUE(args[0]->Is<Constant>());
+    auto val = args[0]->As<Constant>()->Value();
     ASSERT_TRUE(val->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, val->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -44,10 +45,11 @@
     Builder b{mod};
     const auto* inst = b.Bitcast(b.ir.types.Get<type::I32>(), b.Constant(4_i));
 
-    ASSERT_EQ(inst->args.Length(), 1u);
-    ASSERT_NE(inst->args[0], nullptr);
-    ASSERT_EQ(inst->args[0]->Usage().Length(), 1u);
-    EXPECT_EQ(inst->args[0]->Usage()[0], inst);
+    const auto args = inst->Args();
+    ASSERT_EQ(args.Length(), 1u);
+    ASSERT_NE(args[0], nullptr);
+    ASSERT_EQ(args[0]->Usage().Length(), 1u);
+    EXPECT_EQ(args[0]->Usage()[0], inst);
 }
 
 }  // namespace
diff --git a/src/tint/ir/block.cc b/src/tint/ir/block.cc
index 2030c19..f5b5bfb 100644
--- a/src/tint/ir/block.cc
+++ b/src/tint/ir/block.cc
@@ -22,4 +22,13 @@
 
 Block::~Block() = default;
 
+void Block::BranchTo(FlowNode* to, utils::VectorRef<Value*> args) {
+    TINT_ASSERT(IR, to);
+    branch_.target = to;
+    branch_.args = args;
+    if (to) {
+        to->AddInboundBranch(this);
+    }
+}
+
 }  // namespace tint::ir
diff --git a/src/tint/ir/block.h b/src/tint/ir/block.h
index 5af32fe..2b01fd0 100644
--- a/src/tint/ir/block.h
+++ b/src/tint/ir/block.h
@@ -15,6 +15,8 @@
 #ifndef SRC_TINT_IR_BLOCK_H_
 #define SRC_TINT_IR_BLOCK_H_
 
+#include <utility>
+
 #include "src/tint/ir/block_param.h"
 #include "src/tint/ir/branch.h"
 #include "src/tint/ir/flow_node.h"
@@ -30,20 +32,48 @@
   public:
     /// Constructor
     Block();
+    Block(const Block&) = delete;
+    Block(Block&&) = delete;
     ~Block() override;
 
-    /// @returns true if this is a dead block. This can happen in the case like a loop merge block
-    /// which is never reached.
-    bool IsDead() const override { return branch.target == nullptr; }
+    Block& operator=(const Block&) = delete;
+    Block& operator=(Block&&) = delete;
 
-    /// The node this block branches too.
-    Branch branch = {};
+    /// Sets the blocks branch target to the given node.
+    /// @param to the node to branch too
+    /// @param args the branch arguments
+    void BranchTo(FlowNode* to, utils::VectorRef<Value*> args = {});
 
-    /// The instructions in the block
-    utils::Vector<const Instruction*, 16> instructions;
+    /// @returns true if this is block has a branch target set
+    bool HasBranchTarget() const override { return branch_.target != nullptr; }
 
-    /// The parameters passed into the block
-    utils::Vector<const BlockParam*, 0> params;
+    /// @return the node this block branches too.
+    const ir::Branch& Branch() const { return branch_; }
+
+    /// Sets the instructions in the block
+    /// @param instructions the instructions to set
+    void SetInstructions(utils::VectorRef<const Instruction*> instructions) {
+        instructions_ = std::move(instructions);
+    }
+
+    /// @returns the instructions in the block
+    utils::VectorRef<const Instruction*> Instructions() const { return instructions_; }
+    /// @returns the instructions in the block
+    utils::Vector<const Instruction*, 16>& Instructions() { return instructions_; }
+
+    /// Sets the params to the block
+    /// @param params the params for the block
+    void SetParams(utils::VectorRef<const BlockParam*> params) { params_ = std::move(params); }
+    /// @returns the params to the block
+    utils::Vector<const BlockParam*, 0>& Params() { return params_; }
+
+    /// @return the parameters passed into the block
+    utils::VectorRef<const BlockParam*> Params() const { return params_; }
+
+  private:
+    ir::Branch branch_ = {};
+    utils::Vector<const Instruction*, 16> instructions_;
+    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
index f014d19..a3a0be5 100644
--- a/src/tint/ir/block_param.cc
+++ b/src/tint/ir/block_param.cc
@@ -18,7 +18,7 @@
 
 namespace tint::ir {
 
-BlockParam::BlockParam(const type::Type* ty) : type(ty) {}
+BlockParam::BlockParam(const type::Type* ty) : type_(ty) {}
 
 BlockParam::~BlockParam() = default;
 
diff --git a/src/tint/ir/block_param.h b/src/tint/ir/block_param.h
index 8ba68a7..036ddbf 100644
--- a/src/tint/ir/block_param.h
+++ b/src/tint/ir/block_param.h
@@ -34,10 +34,11 @@
     BlockParam& operator=(BlockParam&& inst) = delete;
 
     /// @returns the type of the var
-    const type::Type* Type() const override { return type; }
+    const type::Type* Type() const override { return type_; }
 
+  private:
     /// the result type of the instruction
-    const type::Type* type = nullptr;
+    const type::Type* type_;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index 983815e..b0b04d6 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -32,7 +32,7 @@
 
         // Everything in the module scope must have been const-eval's, so everything will go into a
         // single block. So, we can create the root terminator for the root-block now.
-        ir.root_block->branch.target = CreateRootTerminator();
+        ir.root_block->BranchTo(CreateRootTerminator());
     }
     return ir.root_block;
 }
@@ -56,11 +56,11 @@
     TINT_ASSERT(IR, return_type);
 
     auto* ir_func = ir.flow_nodes.Create<Function>(name, return_type, stage, wg_size);
-    ir_func->start_target = CreateBlock();
-    ir_func->end_target = CreateFunctionTerminator();
+    ir_func->SetStartTarget(CreateBlock());
+    ir_func->SetEndTarget(CreateFunctionTerminator());
 
-    // Function is always branching into the start target
-    ir_func->start_target->inbound_branches.Push(ir_func);
+    // Function is always branching into the Start().target
+    ir_func->StartTarget()->AddInboundBranch(ir_func);
 
     return ir_func;
 }
@@ -69,53 +69,48 @@
     TINT_ASSERT(IR, condition);
 
     auto* ir_if = ir.flow_nodes.Create<If>(condition);
-    ir_if->true_.target = CreateBlock();
-    ir_if->false_.target = CreateBlock();
-    ir_if->merge.target = CreateBlock();
+    ir_if->True().target = CreateBlock();
+    ir_if->False().target = CreateBlock();
+    ir_if->Merge().target = CreateBlock();
 
     // An if always branches to both the true and false block.
-    ir_if->true_.target->inbound_branches.Push(ir_if);
-    ir_if->false_.target->inbound_branches.Push(ir_if);
+    ir_if->True().target->AddInboundBranch(ir_if);
+    ir_if->False().target->AddInboundBranch(ir_if);
 
     return ir_if;
 }
 
 Loop* Builder::CreateLoop() {
     auto* ir_loop = ir.flow_nodes.Create<Loop>();
-    ir_loop->start.target = CreateBlock();
-    ir_loop->continuing.target = CreateBlock();
-    ir_loop->merge.target = CreateBlock();
+    ir_loop->Start().target = CreateBlock();
+    ir_loop->Continuing().target = CreateBlock();
+    ir_loop->Merge().target = CreateBlock();
 
     // A loop always branches to the start block.
-    ir_loop->start.target->inbound_branches.Push(ir_loop);
+    ir_loop->Start().target->AddInboundBranch(ir_loop);
 
     return ir_loop;
 }
 
 Switch* Builder::CreateSwitch(Value* condition) {
     auto* ir_switch = ir.flow_nodes.Create<Switch>(condition);
-    ir_switch->merge.target = CreateBlock();
+    ir_switch->Merge().target = CreateBlock();
     return ir_switch;
 }
 
 Block* Builder::CreateCase(Switch* s, utils::VectorRef<Switch::CaseSelector> selectors) {
-    s->cases.Push(Switch::Case{selectors, {CreateBlock(), utils::Empty}});
+    s->Cases().Push(Switch::Case{selectors, {CreateBlock(), utils::Empty}});
 
-    Block* b = s->cases.Back().start.target->As<Block>();
+    Block* b = s->Cases().Back().Start().target->As<Block>();
     // Switch branches into the case block
-    b->inbound_branches.Push(s);
+    b->AddInboundBranch(s);
     return b;
 }
 
-void Builder::Branch(Block* from, FlowNode* to, utils::VectorRef<Value*> args) {
-    TINT_ASSERT(IR, from);
-    TINT_ASSERT(IR, to);
-    from->branch.target = to;
-    from->branch.args = args;
-    to->inbound_branches.Push(from);
-}
-
-Binary* Builder::CreateBinary(Binary::Kind kind, const type::Type* type, Value* lhs, Value* rhs) {
+Binary* Builder::CreateBinary(enum Binary::Kind kind,
+                              const type::Type* type,
+                              Value* lhs,
+                              Value* rhs) {
     return ir.values.Create<ir::Binary>(kind, type, lhs, rhs);
 }
 
@@ -183,7 +178,7 @@
     return CreateBinary(Binary::Kind::kModulo, type, lhs, rhs);
 }
 
-Unary* Builder::CreateUnary(Unary::Kind kind, const type::Type* type, Value* val) {
+Unary* Builder::CreateUnary(enum Unary::Kind kind, const type::Type* type, Value* val) {
     return ir.values.Create<ir::Unary>(kind, type, val);
 }
 
@@ -247,4 +242,8 @@
     return ir.values.Create<ir::BlockParam>(type);
 }
 
+ir::FunctionParam* Builder::FunctionParam(const type::Type* type) {
+    return ir.values.Create<ir::FunctionParam>(type);
+}
+
 }  // namespace tint::ir
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index 7c0cab6..ebd7a87 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -20,12 +20,14 @@
 #include "src/tint/constant/scalar.h"
 #include "src/tint/ir/binary.h"
 #include "src/tint/ir/bitcast.h"
+#include "src/tint/ir/block_param.h"
 #include "src/tint/ir/builtin.h"
 #include "src/tint/ir/constant.h"
 #include "src/tint/ir/construct.h"
 #include "src/tint/ir/convert.h"
 #include "src/tint/ir/discard.h"
 #include "src/tint/ir/function.h"
+#include "src/tint/ir/function_param.h"
 #include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/load.h"
@@ -96,26 +98,20 @@
     /// @returns the start block for the case flow node
     Block* CreateCase(Switch* s, utils::VectorRef<Switch::CaseSelector> selectors);
 
-    /// Branches the given block to the given flow node.
-    /// @param from the block to branch from
-    /// @param to the node to branch too
-    /// @param args arguments to the branch
-    void Branch(Block* from, FlowNode* to, utils::VectorRef<Value*> args = {});
-
     /// Creates a constant::Value
     /// @param args the arguments
     /// @returns the new constant value
     template <typename T, typename... ARGS>
     utils::traits::EnableIf<utils::traits::IsTypeOrDerived<T, constant::Value>, const T>* create(
         ARGS&&... args) {
-        return ir.constants.Create<T>(std::forward<ARGS>(args)...);
+        return ir.constants_arena.Create<T>(std::forward<ARGS>(args)...);
     }
 
     /// Creates a new ir::Constant
     /// @param val the constant value
     /// @returns the new constant
     ir::Constant* Constant(const constant::Value* val) {
-        return ir.values.Create<ir::Constant>(val);
+        return ir.constants.GetOrCreate(val, [&]() { return ir.values.Create<ir::Constant>(val); });
     }
 
     /// Creates a ir::Constant for an i32 Scalar
@@ -159,7 +155,7 @@
     /// @param lhs the left-hand-side of the operation
     /// @param rhs the right-hand-side of the operation
     /// @returns the operation
-    Binary* CreateBinary(Binary::Kind kind, const type::Type* type, Value* lhs, Value* rhs);
+    Binary* CreateBinary(enum Binary::Kind kind, const type::Type* type, Value* lhs, Value* rhs);
 
     /// Creates an And operation
     /// @param type the result type of the expression
@@ -278,7 +274,7 @@
     /// @param type the result type of the binary expression
     /// @param val the value of the operation
     /// @returns the operation
-    Unary* CreateUnary(Unary::Kind kind, const type::Type* type, Value* val);
+    Unary* CreateUnary(enum Unary::Kind kind, const type::Type* type, Value* val);
 
     /// Creates a Complement operation
     /// @param type the result type of the expression
@@ -360,6 +356,11 @@
     /// @returns the value
     ir::BlockParam* BlockParam(const type::Type* type);
 
+    /// Creates a new `FunctionParam`
+    /// @param type the parameter type
+    /// @returns the value
+    ir::FunctionParam* FunctionParam(const type::Type* type);
+
     /// Retrieves the root block for the module, creating if necessary
     /// @returns the root block
     ir::Block* CreateRootBlockIfNeeded();
diff --git a/src/tint/ir/call.cc b/src/tint/ir/call.cc
index dd503ea..909079f 100644
--- a/src/tint/ir/call.cc
+++ b/src/tint/ir/call.cc
@@ -21,8 +21,8 @@
 namespace tint::ir {
 
 Call::Call(const type::Type* res_ty, utils::VectorRef<Value*> arguments)
-    : result_type(res_ty), args(std::move(arguments)) {
-    for (auto* arg : args) {
+    : result_type_(res_ty), args_(std::move(arguments)) {
+    for (auto* arg : args_) {
         arg->AddUsage(this);
     }
 }
diff --git a/src/tint/ir/call.h b/src/tint/ir/call.h
index 5393810..f4e12f9 100644
--- a/src/tint/ir/call.h
+++ b/src/tint/ir/call.h
@@ -31,13 +31,10 @@
     Call& operator=(Call&& inst) = delete;
 
     /// @returns the type of the value
-    const type::Type* Type() const override { return result_type; }
+    const type::Type* Type() const override { return result_type_; }
 
-    /// The instruction type
-    const type::Type* result_type = nullptr;
-
-    /// The constructor arguments
-    utils::Vector<Value*, 1> args;
+    /// @returns the call arguments
+    utils::VectorRef<Value*> Args() const { return args_; }
 
   protected:
     /// Constructor
@@ -46,6 +43,10 @@
     /// @param result_type the result type
     /// @param args the constructor arguments
     Call(const type::Type* result_type, utils::VectorRef<Value*> args);
+
+  private:
+    const type::Type* result_type_;
+    utils::Vector<Value*, 1> args_;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/constant.cc b/src/tint/ir/constant.cc
index 8b5260c..a49d140 100644
--- a/src/tint/ir/constant.cc
+++ b/src/tint/ir/constant.cc
@@ -18,7 +18,7 @@
 
 namespace tint::ir {
 
-Constant::Constant(const constant::Value* val) : value(val) {}
+Constant::Constant(const constant::Value* val) : value_(val) {}
 
 Constant::~Constant() = default;
 
diff --git a/src/tint/ir/constant.h b/src/tint/ir/constant.h
index 68e0dc7..dcf3e5b 100644
--- a/src/tint/ir/constant.h
+++ b/src/tint/ir/constant.h
@@ -26,13 +26,21 @@
     /// Constructor
     /// @param val the value stored in the constant
     explicit Constant(const constant::Value* val);
+    Constant(const Constant&) = delete;
+    Constant(Constant&&) = delete;
     ~Constant() override;
 
-    /// @returns the type of the constant
-    const type::Type* Type() const override { return value->Type(); }
+    Constant& operator=(const Constant&) = delete;
+    Constant& operator=(Constant&&) = delete;
 
-    /// The constants value
-    const constant::Value* const value;
+    /// @returns the constants value
+    const constant::Value* Value() const { return value_; }
+
+    /// @returns the type of the constant
+    const type::Type* Type() const override { return value_->Type(); }
+
+  private:
+    const constant::Value* const value_;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/constant_test.cc b/src/tint/ir/constant_test.cc
index 7005751..7c3024a 100644
--- a/src/tint/ir/constant_test.cc
+++ b/src/tint/ir/constant_test.cc
@@ -30,13 +30,13 @@
     utils::StringStream str;
 
     auto* c = b.Constant(1.2_f);
-    EXPECT_EQ(1.2_f, c->value->As<constant::Scalar<f32>>()->ValueAs<f32>());
+    EXPECT_EQ(1.2_f, c->Value()->As<constant::Scalar<f32>>()->ValueAs<f32>());
 
-    EXPECT_TRUE(c->value->Is<constant::Scalar<f32>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<f16>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<i32>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<u32>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<bool>>());
+    EXPECT_TRUE(c->Value()->Is<constant::Scalar<f32>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<f16>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<i32>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<u32>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<bool>>());
 }
 
 TEST_F(IR_ConstantTest, f16) {
@@ -46,13 +46,13 @@
     utils::StringStream str;
 
     auto* c = b.Constant(1.1_h);
-    EXPECT_EQ(1.1_h, c->value->As<constant::Scalar<f16>>()->ValueAs<f16>());
+    EXPECT_EQ(1.1_h, c->Value()->As<constant::Scalar<f16>>()->ValueAs<f16>());
 
-    EXPECT_FALSE(c->value->Is<constant::Scalar<f32>>());
-    EXPECT_TRUE(c->value->Is<constant::Scalar<f16>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<i32>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<u32>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<bool>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<f32>>());
+    EXPECT_TRUE(c->Value()->Is<constant::Scalar<f16>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<i32>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<u32>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<bool>>());
 }
 
 TEST_F(IR_ConstantTest, i32) {
@@ -62,13 +62,13 @@
     utils::StringStream str;
 
     auto* c = b.Constant(1_i);
-    EXPECT_EQ(1_i, c->value->As<constant::Scalar<i32>>()->ValueAs<i32>());
+    EXPECT_EQ(1_i, c->Value()->As<constant::Scalar<i32>>()->ValueAs<i32>());
 
-    EXPECT_FALSE(c->value->Is<constant::Scalar<f32>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<f16>>());
-    EXPECT_TRUE(c->value->Is<constant::Scalar<i32>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<u32>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<bool>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<f32>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<f16>>());
+    EXPECT_TRUE(c->Value()->Is<constant::Scalar<i32>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<u32>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<bool>>());
 }
 
 TEST_F(IR_ConstantTest, u32) {
@@ -78,13 +78,13 @@
     utils::StringStream str;
 
     auto* c = b.Constant(2_u);
-    EXPECT_EQ(2_u, c->value->As<constant::Scalar<u32>>()->ValueAs<u32>());
+    EXPECT_EQ(2_u, c->Value()->As<constant::Scalar<u32>>()->ValueAs<u32>());
 
-    EXPECT_FALSE(c->value->Is<constant::Scalar<f32>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<f16>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<i32>>());
-    EXPECT_TRUE(c->value->Is<constant::Scalar<u32>>());
-    EXPECT_FALSE(c->value->Is<constant::Scalar<bool>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<f32>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<f16>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<i32>>());
+    EXPECT_TRUE(c->Value()->Is<constant::Scalar<u32>>());
+    EXPECT_FALSE(c->Value()->Is<constant::Scalar<bool>>());
 }
 
 TEST_F(IR_ConstantTest, bool) {
@@ -95,19 +95,19 @@
         utils::StringStream str;
 
         auto* c = b.Constant(false);
-        EXPECT_FALSE(c->value->As<constant::Scalar<bool>>()->ValueAs<bool>());
+        EXPECT_FALSE(c->Value()->As<constant::Scalar<bool>>()->ValueAs<bool>());
     }
 
     {
         utils::StringStream str;
         auto c = b.Constant(true);
-        EXPECT_TRUE(c->value->As<constant::Scalar<bool>>()->ValueAs<bool>());
+        EXPECT_TRUE(c->Value()->As<constant::Scalar<bool>>()->ValueAs<bool>());
 
-        EXPECT_FALSE(c->value->Is<constant::Scalar<f32>>());
-        EXPECT_FALSE(c->value->Is<constant::Scalar<f16>>());
-        EXPECT_FALSE(c->value->Is<constant::Scalar<i32>>());
-        EXPECT_FALSE(c->value->Is<constant::Scalar<u32>>());
-        EXPECT_TRUE(c->value->Is<constant::Scalar<bool>>());
+        EXPECT_FALSE(c->Value()->Is<constant::Scalar<f32>>());
+        EXPECT_FALSE(c->Value()->Is<constant::Scalar<f16>>());
+        EXPECT_FALSE(c->Value()->Is<constant::Scalar<i32>>());
+        EXPECT_FALSE(c->Value()->Is<constant::Scalar<u32>>());
+        EXPECT_TRUE(c->Value()->Is<constant::Scalar<bool>>());
     }
 }
 
diff --git a/src/tint/ir/debug.cc b/src/tint/ir/debug.cc
index 655c455..85666ff 100644
--- a/src/tint/ir/debug.cc
+++ b/src/tint/ir/debug.cc
@@ -60,81 +60,81 @@
                 if (node_to_name.count(b) == 0) {
                     out << name_for(b) << R"( [label="block"])" << std::endl;
                 }
-                out << name_for(b) << " -> " << name_for(b->branch.target);
+                out << name_for(b) << " -> " << name_for(b->Branch().target);
 
                 // Dashed lines to merge blocks
-                if (merge_nodes.count(b->branch.target) != 0) {
+                if (merge_nodes.count(b->Branch().target) != 0) {
                     out << " [style=dashed]";
                 }
 
                 out << std::endl;
-                Graph(b->branch.target);
+                Graph(b->Branch().target);
             },
             [&](const ir::Switch* s) {
                 out << name_for(s) << R"( [label="switch"])" << std::endl;
-                out << name_for(s->merge.target) << R"( [label="switch merge"])" << std::endl;
-                merge_nodes.insert(s->merge.target);
+                out << name_for(s->Merge().target) << R"( [label="switch merge"])" << std::endl;
+                merge_nodes.insert(s->Merge().target);
 
                 size_t i = 0;
-                for (const auto& c : s->cases) {
-                    out << name_for(c.start.target)
+                for (const auto& c : s->Cases()) {
+                    out << name_for(c.Start().target)
                         << R"( [label="case )" + std::to_string(i++) + R"("])" << std::endl;
                 }
                 out << name_for(s) << " -> {";
-                for (const auto& c : s->cases) {
-                    if (&c != &(s->cases[0])) {
+                for (const auto& c : s->Cases()) {
+                    if (&c != &(s->Cases().Front())) {
                         out << ", ";
                     }
-                    out << name_for(c.start.target);
+                    out << name_for(c.Start().target);
                 }
                 out << "}" << std::endl;
 
-                for (const auto& c : s->cases) {
-                    Graph(c.start.target);
+                for (const auto& c : s->Cases()) {
+                    Graph(c.Start().target);
                 }
-                Graph(s->merge.target);
+                Graph(s->Merge().target);
             },
             [&](const ir::If* i) {
                 out << name_for(i) << R"( [label="if"])" << std::endl;
-                out << name_for(i->true_.target) << R"( [label="true"])" << std::endl;
-                out << name_for(i->false_.target) << R"( [label="false"])" << std::endl;
-                out << name_for(i->merge.target) << R"( [label="if merge"])" << std::endl;
-                merge_nodes.insert(i->merge.target);
+                out << name_for(i->True().target) << R"( [label="true"])" << std::endl;
+                out << name_for(i->False().target) << R"( [label="false"])" << std::endl;
+                out << name_for(i->Merge().target) << R"( [label="if merge"])" << std::endl;
+                merge_nodes.insert(i->Merge().target);
 
                 out << name_for(i) << " -> {";
-                out << name_for(i->true_.target) << ", " << name_for(i->false_.target);
+                out << name_for(i->True().target) << ", " << name_for(i->False().target);
                 out << "}" << std::endl;
 
                 // Subgraph if true/false branches so they draw on the same line
                 out << "subgraph sub_" << name_for(i) << " {" << std::endl;
                 out << R"(rank="same")" << std::endl;
-                out << name_for(i->true_.target) << std::endl;
-                out << name_for(i->false_.target) << std::endl;
+                out << name_for(i->True().target) << std::endl;
+                out << name_for(i->False().target) << std::endl;
                 out << "}" << std::endl;
 
-                Graph(i->true_.target);
-                Graph(i->false_.target);
-                Graph(i->merge.target);
+                Graph(i->True().target);
+                Graph(i->False().target);
+                Graph(i->Merge().target);
             },
             [&](const ir::Loop* l) {
                 out << name_for(l) << R"( [label="loop"])" << std::endl;
-                out << name_for(l->start.target) << R"( [label="start"])" << std::endl;
-                out << name_for(l->continuing.target) << R"( [label="continuing"])" << std::endl;
-                out << name_for(l->merge.target) << R"( [label="loop merge"])" << std::endl;
-                merge_nodes.insert(l->merge.target);
+                out << name_for(l->Start().target) << R"( [label="start"])" << std::endl;
+                out << name_for(l->Continuing().target) << R"( [label="continuing"])" << std::endl;
+                out << name_for(l->Merge().target) << R"( [label="loop merge"])" << std::endl;
+                merge_nodes.insert(l->Merge().target);
 
                 // Subgraph the continuing and merge so they get drawn on the same line
                 out << "subgraph sub_" << name_for(l) << " {" << std::endl;
                 out << R"(rank="same")" << std::endl;
-                out << name_for(l->continuing.target) << std::endl;
-                out << name_for(l->merge.target) << std::endl;
+                out << name_for(l->Continuing().target) << std::endl;
+                out << name_for(l->Merge().target) << std::endl;
                 out << "}" << std::endl;
 
-                out << name_for(l) << " -> " << name_for(l->start.target) << std::endl;
+                out << name_for(l) << " -> " << name_for(l->Start().target) << std::endl;
 
-                Graph(l->start.target);
-                Graph(l->continuing.target);
-                Graph(l->merge.target);
+                Graph(l->Start().target);
+                Graph(l->Continuing().target);
+                Graph(l->Merge().target);
             },
             [&](const ir::FunctionTerminator*) {
                 // Already done
@@ -145,10 +145,10 @@
     for (const auto* func : mod->functions) {
         // Cluster each function to label and draw a box around it.
         out << "subgraph cluster_" << name_for(func) << " {" << std::endl;
-        out << R"(label=")" << func->name.Name() << R"(")" << std::endl;
-        out << name_for(func->start_target) << R"( [label="start"])" << std::endl;
-        out << name_for(func->end_target) << R"( [label="end"])" << std::endl;
-        Graph(func->start_target);
+        out << R"(label=")" << func->Name().Name() << R"(")" << std::endl;
+        out << name_for(func->StartTarget()) << R"( [label="start"])" << std::endl;
+        out << name_for(func->EndTarget()) << R"( [label="end"])" << std::endl;
+        Graph(func->StartTarget());
         out << "}" << std::endl;
     }
     out << "}";
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index f68b21a..76af740 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -81,7 +81,7 @@
 }
 
 void Disassembler::EmitBlockInstructions(const Block* b) {
-    for (const auto* inst : b->instructions) {
+    for (const auto* inst : b->Instructions()) {
         Indent();
         EmitInstruction(inst);
         out_ << std::endl;
@@ -114,25 +114,31 @@
         [&](const ir::Function* f) {
             TINT_SCOPED_ASSIGNMENT(in_function_, true);
 
-            Indent() << "%fn" << IdOf(f) << " = func " << f->name.Name()
-                     << "():" << f->return_type->FriendlyName();
+            Indent() << "%fn" << IdOf(f) << " = func " << f->Name().Name() << "(";
+            for (auto* p : f->Params()) {
+                if (p != f->Params().Front()) {
+                    out_ << ", ";
+                }
+                out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
+            }
+            out_ << "):" << f->ReturnType()->FriendlyName();
 
-            if (f->pipeline_stage != Function::PipelineStage::kUndefined) {
-                out_ << " [@" << f->pipeline_stage;
+            if (f->Stage() != Function::PipelineStage::kUndefined) {
+                out_ << " [@" << f->Stage();
 
-                if (f->workgroup_size) {
-                    auto arr = f->workgroup_size.value();
+                if (f->WorkgroupSize()) {
+                    auto arr = f->WorkgroupSize().value();
                     out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2]
                          << ")";
                 }
 
-                if (!f->return_attributes.IsEmpty()) {
+                if (!f->ReturnAttributes().IsEmpty()) {
                     out_ << " ra:";
 
-                    for (auto attr : f->return_attributes) {
+                    for (auto attr : f->ReturnAttributes()) {
                         out_ << " @" << attr;
                         if (attr == Function::ReturnAttribute::kLocation) {
-                            out_ << "(" << f->return_location.value() << ")";
+                            out_ << "(" << f->ReturnLocation().value() << ")";
                         }
                     }
                 }
@@ -143,23 +149,23 @@
 
             {
                 ScopedIndent func_indent(indent_size_);
-                ScopedStopNode scope(stop_nodes_, f->end_target);
-                Walk(f->start_target);
+                ScopedStopNode scope(stop_nodes_, f->EndTarget());
+                Walk(f->StartTarget());
             }
             out_ << "} ";
-            Walk(f->end_target);
+            Walk(f->EndTarget());
         },
         [&](const ir::Block* b) {
             // If this block is dead, nothing to do
-            if (b->IsDead()) {
+            if (!b->HasBranchTarget()) {
                 return;
             }
 
             Indent() << "%fn" << IdOf(b) << " = block";
-            if (!b->params.IsEmpty()) {
+            if (!b->Params().IsEmpty()) {
                 out_ << " (";
-                for (auto* p : b->params) {
-                    if (p != b->params.Front()) {
+                for (const auto* p : b->Params()) {
+                    if (p != b->Params().Front()) {
                         out_ << ", ";
                     }
                     EmitValue(p);
@@ -175,20 +181,20 @@
             Indent() << "}";
 
             std::string suffix = "";
-            if (b->branch.target->Is<FunctionTerminator>()) {
+            if (b->Branch().target->Is<FunctionTerminator>()) {
                 out_ << " -> %func_end";
                 suffix = "return";
-            } else if (b->branch.target->Is<RootTerminator>()) {
+            } else if (b->Branch().target->Is<RootTerminator>()) {
                 // Nothing to do
             } else {
                 out_ << " -> "
-                     << "%fn" << IdOf(b->branch.target);
+                     << "%fn" << IdOf(b->Branch().target);
                 suffix = "branch";
             }
-            if (!b->branch.args.IsEmpty()) {
+            if (!b->Branch().args.IsEmpty()) {
                 out_ << " ";
-                for (const auto* v : b->branch.args) {
-                    if (v != b->branch.args.Front()) {
+                for (const auto* v : b->Branch().args) {
+                    if (v != b->Branch().args.Front()) {
                         out_ << ", ";
                     }
                     EmitValue(v);
@@ -199,18 +205,18 @@
             }
             out_ << std::endl;
 
-            if (!b->branch.target->Is<FunctionTerminator>()) {
+            if (!b->Branch().target->Is<FunctionTerminator>()) {
                 out_ << std::endl;
             }
 
-            Walk(b->branch.target);
+            Walk(b->Branch().target);
         },
         [&](const ir::Switch* s) {
             Indent() << "%fn" << IdOf(s) << " = switch ";
-            EmitValue(s->condition);
+            EmitValue(s->Condition());
             out_ << " [";
-            for (const auto& c : s->cases) {
-                if (&c != &s->cases.Front()) {
+            for (const auto& c : s->Cases()) {
+                if (&c != &s->Cases().Front()) {
                     out_ << ", ";
                 }
                 out_ << "c: (";
@@ -225,17 +231,17 @@
                         EmitValue(selector.val);
                     }
                 }
-                out_ << ", %fn" << IdOf(c.start.target) << ")";
+                out_ << ", %fn" << IdOf(c.Start().target) << ")";
             }
-            if (s->merge.target->IsConnected()) {
-                out_ << ", m: %fn" << IdOf(s->merge.target);
+            if (s->Merge().target->IsConnected()) {
+                out_ << ", m: %fn" << IdOf(s->Merge().target);
             }
             out_ << "]" << std::endl;
 
             {
                 ScopedIndent switch_indent(indent_size_);
-                ScopedStopNode scope(stop_nodes_, s->merge.target);
-                for (const auto& c : s->cases) {
+                ScopedStopNode scope(stop_nodes_, s->Merge().target);
+                for (const auto& c : s->Cases()) {
                     Indent() << "# case ";
                     for (const auto& selector : c.selectors) {
                         if (&selector != &c.selectors.Front()) {
@@ -249,86 +255,86 @@
                         }
                     }
                     out_ << std::endl;
-                    Walk(c.start.target);
+                    Walk(c.Start().target);
                 }
             }
 
-            if (s->merge.target->IsConnected()) {
+            if (s->Merge().target->IsConnected()) {
                 Indent() << "# switch merge" << std::endl;
-                Walk(s->merge.target);
+                Walk(s->Merge().target);
             }
         },
         [&](const ir::If* i) {
             Indent() << "%fn" << IdOf(i) << " = if ";
-            EmitValue(i->condition);
+            EmitValue(i->Condition());
 
-            bool has_true = !i->true_.target->IsDead();
-            bool has_false = !i->false_.target->IsDead();
+            bool has_true = i->True().target->HasBranchTarget();
+            bool has_false = i->False().target->HasBranchTarget();
 
             out_ << " [";
             if (has_true) {
-                out_ << "t: %fn" << IdOf(i->true_.target);
+                out_ << "t: %fn" << IdOf(i->True().target);
             }
             if (has_false) {
                 if (has_true) {
                     out_ << ", ";
                 }
-                out_ << "f: %fn" << IdOf(i->false_.target);
+                out_ << "f: %fn" << IdOf(i->False().target);
             }
-            if (i->merge.target->IsConnected()) {
-                out_ << ", m: %fn" << IdOf(i->merge.target);
+            if (i->Merge().target->IsConnected()) {
+                out_ << ", m: %fn" << IdOf(i->Merge().target);
             }
             out_ << "]" << std::endl;
 
             {
                 ScopedIndent if_indent(indent_size_);
-                ScopedStopNode scope(stop_nodes_, i->merge.target);
+                ScopedStopNode scope(stop_nodes_, i->Merge().target);
 
                 if (has_true) {
                     Indent() << "# true branch" << std::endl;
-                    Walk(i->true_.target);
+                    Walk(i->True().target);
                 }
 
                 if (has_false) {
                     Indent() << "# false branch" << std::endl;
-                    Walk(i->false_.target);
+                    Walk(i->False().target);
                 }
             }
 
-            if (i->merge.target->IsConnected()) {
+            if (i->Merge().target->IsConnected()) {
                 Indent() << "# if merge" << std::endl;
-                Walk(i->merge.target);
+                Walk(i->Merge().target);
             }
         },
         [&](const ir::Loop* l) {
-            Indent() << "%fn" << IdOf(l) << " = loop [s: %fn" << IdOf(l->start.target);
+            Indent() << "%fn" << IdOf(l) << " = loop [s: %fn" << IdOf(l->Start().target);
 
-            if (l->continuing.target->IsConnected()) {
-                out_ << ", c: %fn" << IdOf(l->continuing.target);
+            if (l->Continuing().target->IsConnected()) {
+                out_ << ", c: %fn" << IdOf(l->Continuing().target);
             }
-            if (l->merge.target->IsConnected()) {
-                out_ << ", m: %fn" << IdOf(l->merge.target);
+            if (l->Merge().target->IsConnected()) {
+                out_ << ", m: %fn" << IdOf(l->Merge().target);
             }
             out_ << "]" << std::endl;
 
             {
-                ScopedStopNode loop_scope(stop_nodes_, l->merge.target);
+                ScopedStopNode loop_scope(stop_nodes_, l->Merge().target);
                 ScopedIndent loop_indent(indent_size_);
                 {
-                    ScopedStopNode inner_scope(stop_nodes_, l->continuing.target);
+                    ScopedStopNode inner_scope(stop_nodes_, l->Continuing().target);
                     Indent() << "# loop start" << std::endl;
-                    Walk(l->start.target);
+                    Walk(l->Start().target);
                 }
 
-                if (l->continuing.target->IsConnected()) {
+                if (l->Continuing().target->IsConnected()) {
                     Indent() << "# loop continuing" << std::endl;
-                    Walk(l->continuing.target);
+                    Walk(l->Continuing().target);
                 }
             }
 
-            if (l->merge.target->IsConnected()) {
+            if (l->Merge().target->IsConnected()) {
                 Indent() << "# loop merge" << std::endl;
-                Walk(l->merge.target);
+                Walk(l->Merge().target);
             }
         },
         [&](const ir::FunctionTerminator*) {
@@ -401,12 +407,13 @@
                         }
                     });
             };
-            emit(constant->value);
+            emit(constant->Value());
         },
         [&](const ir::Instruction* i) { out_ << "%" << IdOf(i); },
         [&](const ir::BlockParam* p) {
             out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
         },
+        [&](const ir::FunctionParam* p) { out_ << "%" << IdOf(p); },
         [&](Default) { out_ << "Unknown value: " << val->TypeInfo().name; });
 }
 
@@ -438,18 +445,18 @@
         [&](const ir::Load* l) {
             EmitValueWithType(l);
             out_ << " = load ";
-            EmitValue(l->from);
+            EmitValue(l->From());
         },
         [&](const ir::Store* s) {
             out_ << "store ";
-            EmitValue(s->to);
+            EmitValue(s->To());
             out_ << ", ";
-            EmitValue(s->from);
+            EmitValue(s->From());
         },
         [&](const ir::UserCall* uc) {
             EmitValueWithType(uc);
-            out_ << " = call " << uc->name.Name();
-            if (uc->args.Length() > 0) {
+            out_ << " = call " << uc->Name().Name();
+            if (!uc->Args().IsEmpty()) {
                 out_ << ", ";
             }
             EmitArgs(uc);
@@ -457,16 +464,16 @@
         [&](const ir::Var* v) {
             EmitValueWithType(v);
             out_ << " = var";
-            if (v->initializer) {
+            if (v->Initializer()) {
                 out_ << ", ";
-                EmitValue(v->initializer);
+                EmitValue(v->Initializer());
             }
         });
 }
 
 void Disassembler::EmitArgs(const Call* call) {
     bool first = true;
-    for (const auto* arg : call->args) {
+    for (const auto* arg : call->Args()) {
         if (!first) {
             out_ << ", ";
         }
@@ -478,7 +485,7 @@
 void Disassembler::EmitBinary(const Binary* b) {
     EmitValueWithType(b);
     out_ << " = ";
-    switch (b->kind) {
+    switch (b->Kind()) {
         case Binary::Kind::kAdd:
             out_ << "add";
             break;
@@ -537,7 +544,7 @@
 void Disassembler::EmitUnary(const Unary* u) {
     EmitValueWithType(u);
     out_ << " = ";
-    switch (u->kind) {
+    switch (u->Kind()) {
         case Unary::Kind::kComplement:
             out_ << "complement";
             break;
diff --git a/src/tint/ir/flow_node.h b/src/tint/ir/flow_node.h
index 905f077..289873b 100644
--- a/src/tint/ir/flow_node.h
+++ b/src/tint/ir/flow_node.h
@@ -25,22 +25,30 @@
   public:
     ~FlowNode() override;
 
-    /// The list of flow nodes which branch into this node. This list maybe empty for several
-    /// reasons:
-    ///   - Node is a start node
-    ///   - Node is a merge target outside control flow (if that returns in both branches)
-    ///   - Node is a continue target outside control flow (loop that returns)
-    utils::Vector<FlowNode*, 2> inbound_branches;
-
     /// @returns true if this node has inbound branches and branches out
-    bool IsConnected() const { return !IsDead() && !inbound_branches.IsEmpty(); }
+    bool IsConnected() const { return HasBranchTarget() && !inbound_branches_.IsEmpty(); }
 
-    /// @returns true if the node does not branch out
-    virtual bool IsDead() const { return false; }
+    /// @returns true if the node has a branch target
+    virtual bool HasBranchTarget() const { return false; }
+
+    /// @returns the inbound branch list for the flow node
+    utils::VectorRef<FlowNode*> InboundBranches() const { return inbound_branches_; }
+
+    /// Adds the given node to the inbound branches
+    /// @param node the node to add
+    void AddInboundBranch(FlowNode* node) { inbound_branches_.Push(node); }
 
   protected:
     /// Constructor
     FlowNode();
+
+  private:
+    /// The list of flow nodes which branch into this node. This list maybe empty for several
+    /// reasons:
+    ///   - Node is a start node
+    ///   - Node is a merge target outside control flow (e.g. an if that returns in both branches)
+    ///   - Node is a continue target outside control flow (e.g. a loop that returns)
+    utils::Vector<FlowNode*, 2> inbound_branches_;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index bb9c3f9..595db58 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -98,17 +98,13 @@
 
 using ResultType = utils::Result<Module, diag::List>;
 
-bool IsBranched(const Block* b) {
-    return b->branch.target != nullptr;
-}
-
 bool IsConnected(const FlowNode* b) {
     // Function is always connected as it's the start.
     if (b->Is<ir::Function>()) {
         return true;
     }
 
-    for (auto* parent : b->inbound_branches) {
+    for (auto* parent : b->InboundBranches()) {
         if (IsConnected(parent)) {
             return true;
         }
@@ -146,7 +142,7 @@
             /* src */ {&program_->Symbols()},
             /* dst */ {&builder_.ir.symbols, &builder_.ir.types},
         },
-        /* dst */ {&builder_.ir.constants},
+        /* dst */ {&builder_.ir.constants_arena},
     };
 
     /// The stack of flow control blocks.
@@ -184,14 +180,14 @@
 
     void BranchTo(FlowNode* node, utils::VectorRef<Value*> args = {}) {
         TINT_ASSERT(IR, current_flow_block_);
-        TINT_ASSERT(IR, !IsBranched(current_flow_block_));
+        TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget());
 
-        builder_.Branch(current_flow_block_, node, args);
+        current_flow_block_->BranchTo(node, args);
         current_flow_block_ = nullptr;
     }
 
     void BranchToIfNeeded(FlowNode* node) {
-        if (!current_flow_block_ || IsBranched(current_flow_block_)) {
+        if (!current_flow_block_ || current_flow_block_->HasBranchTarget()) {
             return;
         }
         BranchTo(node);
@@ -271,20 +267,17 @@
         if (ast_func->IsEntryPoint()) {
             switch (ast_func->PipelineStage()) {
                 case ast::PipelineStage::kVertex:
-                    ir_func->pipeline_stage = Function::PipelineStage::kVertex;
+                    ir_func->SetStage(Function::PipelineStage::kVertex);
                     break;
                 case ast::PipelineStage::kFragment:
-                    ir_func->pipeline_stage = Function::PipelineStage::kFragment;
+                    ir_func->SetStage(Function::PipelineStage::kFragment);
                     break;
                 case ast::PipelineStage::kCompute: {
-                    ir_func->pipeline_stage = Function::PipelineStage::kCompute;
+                    ir_func->SetStage(Function::PipelineStage::kCompute);
 
                     auto wg_size = sem->WorkgroupSize();
-                    ir_func->workgroup_size = {
-                        wg_size[0].value(),
-                        wg_size[1].value_or(1),
-                        wg_size[2].value_or(1),
-                    };
+                    ir_func->SetWorkgroupSize(wg_size[0].value(), wg_size[1].value_or(1),
+                                              wg_size[2].value_or(1));
                     break;
                 }
                 default: {
@@ -293,14 +286,15 @@
                 }
             }
 
+            utils::Vector<Function::ReturnAttribute, 1> return_attributes;
             for (auto* attr : ast_func->return_type_attributes) {
                 tint::Switch(
                     attr,  //
                     [&](const ast::LocationAttribute*) {
-                        ir_func->return_attributes.Push(Function::ReturnAttribute::kLocation);
+                        return_attributes.Push(Function::ReturnAttribute::kLocation);
                     },
                     [&](const ast::InvariantAttribute*) {
-                        ir_func->return_attributes.Push(Function::ReturnAttribute::kInvariant);
+                        return_attributes.Push(Function::ReturnAttribute::kInvariant);
                     },
                     [&](const ast::BuiltinAttribute* b) {
                         if (auto* ident_sem =
@@ -309,16 +303,13 @@
                                     ->As<sem::BuiltinEnumExpression<builtin::BuiltinValue>>()) {
                             switch (ident_sem->Value()) {
                                 case builtin::BuiltinValue::kPosition:
-                                    ir_func->return_attributes.Push(
-                                        Function::ReturnAttribute::kPosition);
+                                    return_attributes.Push(Function::ReturnAttribute::kPosition);
                                     break;
                                 case builtin::BuiltinValue::kFragDepth:
-                                    ir_func->return_attributes.Push(
-                                        Function::ReturnAttribute::kFragDepth);
+                                    return_attributes.Push(Function::ReturnAttribute::kFragDepth);
                                     break;
                                 case builtin::BuiltinValue::kSampleMask:
-                                    ir_func->return_attributes.Push(
-                                        Function::ReturnAttribute::kSampleMask);
+                                    return_attributes.Push(Function::ReturnAttribute::kSampleMask);
                                     break;
                                 default:
                                     TINT_ICE(IR, diagnostics_)
@@ -332,20 +323,34 @@
                         }
                     });
             }
+            ir_func->SetReturnAttributes(return_attributes);
         }
-        ir_func->return_location = sem->ReturnLocation();
+        ir_func->SetReturnLocation(sem->ReturnLocation());
+
+        scopes_.Push();
+        TINT_DEFER(scopes_.Pop());
+
+        utils::Vector<FunctionParam*, 1> params;
+        for (auto* p : ast_func->params) {
+            const auto* param_sem = program_->Sem().Get(p);
+            auto* ty = param_sem->Type()->Clone(clone_ctx_.type_ctx);
+            auto* param = builder_.FunctionParam(ty);
+
+            scopes_.Set(p->name->symbol, param);
+            builder_.ir.SetName(param, p->name->symbol.NameView());
+            params.Push(param);
+        }
+        ir_func->SetParams(params);
 
         {
             FlowStackScope scope(this, ir_func);
 
-            current_flow_block_ = ir_func->start_target;
+            current_flow_block_ = ir_func->StartTarget();
             EmitBlock(ast_func->body);
 
-            // TODO(dsinclair): Store parameters
-
             // If the branch target has already been set then a `return` was called. Only set in the
             // case where `return` wasn't called.
-            BranchToIfNeeded(current_function_->end_target);
+            BranchToIfNeeded(current_function_->EndTarget());
         }
 
         TINT_ASSERT(IR, flow_stack_.IsEmpty());
@@ -359,7 +364,7 @@
 
             // If the current flow block has a branch target then the rest of the statements in this
             // block are dead code. Skip them.
-            if (!current_flow_block_ || IsBranched(current_flow_block_)) {
+            if (!current_flow_block_ || current_flow_block_->HasBranchTarget()) {
                 break;
             }
         }
@@ -414,7 +419,7 @@
             return;
         }
         auto store = builder_.Store(lhs.Get(), rhs.Get());
-        current_flow_block_->instructions.Push(store);
+        current_flow_block_->Instructions().Push(store);
     }
 
     void EmitIncrementDecrement(const ast::IncrementDecrementStatement* stmt) {
@@ -425,7 +430,7 @@
 
         // Load from the LHS.
         auto* lhs_value = builder_.Load(lhs.Get());
-        current_flow_block_->instructions.Push(lhs_value);
+        current_flow_block_->Instructions().Push(lhs_value);
 
         auto* ty = lhs_value->Type();
 
@@ -438,10 +443,10 @@
         } else {
             inst = builder_.Subtract(ty, lhs_value, rhs);
         }
-        current_flow_block_->instructions.Push(inst);
+        current_flow_block_->Instructions().Push(inst);
 
         auto store = builder_.Store(lhs.Get(), inst);
-        current_flow_block_->instructions.Push(store);
+        current_flow_block_->Instructions().Push(store);
     }
 
     void EmitCompoundAssignment(const ast::CompoundAssignmentStatement* stmt) {
@@ -457,7 +462,7 @@
 
         // Load from the LHS.
         auto* lhs_value = builder_.Load(lhs.Get());
-        current_flow_block_->instructions.Push(lhs_value);
+        current_flow_block_->Instructions().Push(lhs_value);
 
         auto* ty = lhs_value->Type();
 
@@ -507,10 +512,10 @@
                 TINT_ICE(IR, diagnostics_) << "missing binary operand type";
                 return;
         }
-        current_flow_block_->instructions.Push(inst);
+        current_flow_block_->Instructions().Push(inst);
 
         auto store = builder_.Store(lhs.Get(), inst);
-        current_flow_block_->instructions.Push(store);
+        current_flow_block_->Instructions().Push(store);
     }
 
     void EmitBlock(const ast::BlockStatement* block) {
@@ -538,27 +543,27 @@
         {
             FlowStackScope scope(this, if_node);
 
-            current_flow_block_ = if_node->true_.target->As<Block>();
+            current_flow_block_ = if_node->True().target->As<Block>();
             EmitBlock(stmt->body);
 
-            // If the true branch did not execute control flow, then go to the merge target
-            BranchToIfNeeded(if_node->merge.target);
+            // If the true branch did not execute control flow, then go to the Merge().target
+            BranchToIfNeeded(if_node->Merge().target);
 
-            current_flow_block_ = if_node->false_.target->As<Block>();
+            current_flow_block_ = if_node->False().target->As<Block>();
             if (stmt->else_statement) {
                 EmitStatement(stmt->else_statement);
             }
 
-            // If the false branch did not execute control flow, then go to the merge target
-            BranchToIfNeeded(if_node->merge.target);
+            // If the false branch did not execute control flow, then go to the Merge().target
+            BranchToIfNeeded(if_node->Merge().target);
         }
         current_flow_block_ = nullptr;
 
         // If both branches went somewhere, then they both returned, continued or broke. So, there
         // is no need for the if merge-block and there is nothing to branch to the merge block
         // anyway.
-        if (IsConnected(if_node->merge.target)) {
-            current_flow_block_ = if_node->merge.target->As<Block>();
+        if (IsConnected(if_node->Merge().target)) {
+            current_flow_block_ = if_node->Merge().target->As<Block>();
         }
     }
 
@@ -572,7 +577,7 @@
         {
             FlowStackScope scope(this, loop_node);
 
-            current_flow_block_ = loop_node->start.target->As<Block>();
+            current_flow_block_ = loop_node->Start().target->As<Block>();
 
             // The loop doesn't use EmitBlock because it needs the scope stack to not get popped
             // until after the continuing block.
@@ -581,21 +586,22 @@
             EmitStatements(stmt->body->statements);
 
             // The current block didn't `break`, `return` or `continue`, go to the continuing block.
-            BranchToIfNeeded(loop_node->continuing.target);
+            BranchToIfNeeded(loop_node->Continuing().target);
 
-            current_flow_block_ = loop_node->continuing.target->As<Block>();
+            current_flow_block_ = loop_node->Continuing().target->As<Block>();
             if (stmt->continuing) {
                 EmitBlock(stmt->continuing);
             }
 
             // Branch back to the start node if the continue target didn't branch out already
-            BranchToIfNeeded(loop_node->start.target);
+            BranchToIfNeeded(loop_node->Start().target);
         }
 
         // The loop merge can get disconnected if the loop returns directly, or the continuing
-        // target branches, eventually, to the merge, but nothing branched to the continuing target.
-        current_flow_block_ = loop_node->merge.target->As<Block>();
-        if (!IsConnected(loop_node->merge.target)) {
+        // target branches, eventually, to the merge, but nothing branched to the
+        // Continuing().target.
+        current_flow_block_ = loop_node->Merge().target->As<Block>();
+        if (!IsConnected(loop_node->Merge().target)) {
             current_flow_block_ = nullptr;
         }
     }
@@ -603,9 +609,8 @@
     void EmitWhile(const ast::WhileStatement* stmt) {
         auto* loop_node = builder_.CreateLoop();
         // Continue is always empty, just go back to the start
-        TINT_ASSERT(IR, loop_node->continuing.target->Is<Block>());
-        builder_.Branch(loop_node->continuing.target->As<Block>(), loop_node->start.target,
-                        utils::Empty);
+        TINT_ASSERT(IR, loop_node->Continuing().target->Is<Block>());
+        loop_node->Continuing().target->As<Block>()->BranchTo(loop_node->Start().target);
 
         BranchTo(loop_node);
 
@@ -614,9 +619,9 @@
         {
             FlowStackScope scope(this, loop_node);
 
-            current_flow_block_ = loop_node->start.target->As<Block>();
+            current_flow_block_ = loop_node->Start().target->As<Block>();
 
-            // Emit the while condition into the start target of the loop
+            // Emit the while condition into the Start().target of the loop
             auto reg = EmitExpression(stmt->condition);
             if (!reg) {
                 return;
@@ -624,31 +629,24 @@
 
             // Create an `if (cond) {} else {break;}` control flow
             auto* if_node = builder_.CreateIf(reg.Get());
-            TINT_ASSERT(IR, if_node->true_.target->Is<Block>());
-            builder_.Branch(if_node->true_.target->As<Block>(), if_node->merge.target,
-                            utils::Empty);
-
-            TINT_ASSERT(IR, if_node->false_.target->Is<Block>());
-            builder_.Branch(if_node->false_.target->As<Block>(), loop_node->merge.target,
-                            utils::Empty);
+            if_node->True().target->As<Block>()->BranchTo(if_node->Merge().target);
+            if_node->False().target->As<Block>()->BranchTo(loop_node->Merge().target);
 
             BranchTo(if_node);
 
-            current_flow_block_ = if_node->merge.target->As<Block>();
+            current_flow_block_ = if_node->Merge().target->As<Block>();
             EmitBlock(stmt->body);
 
-            BranchToIfNeeded(loop_node->continuing.target);
+            BranchToIfNeeded(loop_node->Continuing().target);
         }
-        // The while loop always has a path to the merge target as the break statement comes before
-        // anything inside the loop.
-        current_flow_block_ = loop_node->merge.target->As<Block>();
+        // The while loop always has a path to the Merge().target as the break statement comes
+        // before anything inside the loop.
+        current_flow_block_ = loop_node->Merge().target->As<Block>();
     }
 
     void EmitForLoop(const ast::ForLoopStatement* stmt) {
         auto* loop_node = builder_.CreateLoop();
-        TINT_ASSERT(IR, loop_node->continuing.target->Is<Block>());
-        builder_.Branch(loop_node->continuing.target->As<Block>(), loop_node->start.target,
-                        utils::Empty);
+        loop_node->Continuing().target->As<Block>()->BranchTo(loop_node->Start().target);
 
         // Make sure the initializer ends up in a contained scope
         scopes_.Push();
@@ -666,7 +664,7 @@
         {
             FlowStackScope scope(this, loop_node);
 
-            current_flow_block_ = loop_node->start.target->As<Block>();
+            current_flow_block_ = loop_node->Start().target->As<Block>();
 
             if (stmt->condition) {
                 // Emit the condition into the target target of the loop
@@ -677,30 +675,25 @@
 
                 // Create an `if (cond) {} else {break;}` control flow
                 auto* if_node = builder_.CreateIf(reg.Get());
-                TINT_ASSERT(IR, if_node->true_.target->Is<Block>());
-                builder_.Branch(if_node->true_.target->As<Block>(), if_node->merge.target,
-                                utils::Empty);
-
-                TINT_ASSERT(IR, if_node->false_.target->Is<Block>());
-                builder_.Branch(if_node->false_.target->As<Block>(), loop_node->merge.target,
-                                utils::Empty);
+                if_node->True().target->As<Block>()->BranchTo(if_node->Merge().target);
+                if_node->False().target->As<Block>()->BranchTo(loop_node->Merge().target);
 
                 BranchTo(if_node);
-                current_flow_block_ = if_node->merge.target->As<Block>();
+                current_flow_block_ = if_node->Merge().target->As<Block>();
             }
 
             EmitBlock(stmt->body);
-            BranchToIfNeeded(loop_node->continuing.target);
+            BranchToIfNeeded(loop_node->Continuing().target);
 
             if (stmt->continuing) {
-                current_flow_block_ = loop_node->continuing.target->As<Block>();
+                current_flow_block_ = loop_node->Continuing().target->As<Block>();
                 EmitStatement(stmt->continuing);
             }
         }
 
-        // The while loop always has a path to the merge target as the break statement comes before
-        // anything inside the loop.
-        current_flow_block_ = loop_node->merge.target->As<Block>();
+        // The while loop always has a path to the Merge().target as the break statement comes
+        // before anything inside the loop.
+        current_flow_block_ = loop_node->Merge().target->As<Block>();
     }
 
     void EmitSwitch(const ast::SwitchStatement* stmt) {
@@ -732,13 +725,13 @@
                 current_flow_block_ = builder_.CreateCase(switch_node, selectors);
                 EmitBlock(c->Body()->Declaration());
 
-                BranchToIfNeeded(switch_node->merge.target);
+                BranchToIfNeeded(switch_node->Merge().target);
             }
         }
         current_flow_block_ = nullptr;
 
-        if (IsConnected(switch_node->merge.target)) {
-            current_flow_block_ = switch_node->merge.target->As<Block>();
+        if (IsConnected(switch_node->Merge().target)) {
+            current_flow_block_ = switch_node->Merge().target->As<Block>();
         }
     }
 
@@ -752,7 +745,7 @@
             ret_value.Push(ret.Get());
         }
 
-        BranchTo(current_function_->end_target, std::move(ret_value));
+        BranchTo(current_function_->EndTarget(), std::move(ret_value));
     }
 
     void EmitBreak(const ast::BreakStatement*) {
@@ -760,9 +753,9 @@
         TINT_ASSERT(IR, current_control);
 
         if (auto* c = current_control->As<Loop>()) {
-            BranchTo(c->merge.target);
+            BranchTo(c->Merge().target);
         } else if (auto* s = current_control->As<Switch>()) {
-            BranchTo(s->merge.target);
+            BranchTo(s->Merge().target);
         } else {
             TINT_UNREACHABLE(IR, diagnostics_);
         }
@@ -773,7 +766,7 @@
         TINT_ASSERT(IR, current_control);
 
         if (auto* c = current_control->As<Loop>()) {
-            BranchTo(c->continuing.target);
+            BranchTo(c->Continuing().target);
         } else {
             TINT_UNREACHABLE(IR, diagnostics_);
         }
@@ -785,7 +778,7 @@
     // figuring out the multi-level exit that is triggered.
     void EmitDiscard(const ast::DiscardStatement*) {
         auto* inst = builder_.Discard();
-        current_flow_block_->instructions.Push(inst);
+        current_flow_block_->Instructions().Push(inst);
     }
 
     void EmitBreakIf(const ast::BreakIfStatement* stmt) {
@@ -806,17 +799,17 @@
 
         auto* loop = current_control->As<Loop>();
 
-        current_flow_block_ = if_node->true_.target->As<Block>();
-        BranchTo(loop->merge.target);
+        current_flow_block_ = if_node->True().target->As<Block>();
+        BranchTo(loop->Merge().target);
 
-        current_flow_block_ = if_node->false_.target->As<Block>();
-        BranchTo(if_node->merge.target);
+        current_flow_block_ = if_node->False().target->As<Block>();
+        BranchTo(if_node->Merge().target);
 
-        current_flow_block_ = if_node->merge.target->As<Block>();
+        current_flow_block_ = if_node->Merge().target->As<Block>();
 
         // The `break-if` has to be the last item in the continuing block. The false branch of the
         // `break-if` will always take us back to the start of the loop.
-        BranchTo(loop->start.target);
+        BranchTo(loop->Start().target);
     }
 
     utils::Result<Value*> EmitExpression(const ast::Expression* expr) {
@@ -863,7 +856,7 @@
         // 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);
+            current_flow_block_->Instructions().Push(load);
             return load;
         }
 
@@ -882,14 +875,14 @@
                     ref->Access());
 
                 auto* val = builder_.Declare(ty);
-                current_flow_block_->instructions.Push(val);
+                current_flow_block_->Instructions().Push(val);
 
                 if (v->initializer) {
                     auto init = EmitExpression(v->initializer);
                     if (!init) {
                         return;
                     }
-                    val->initializer = init.Get();
+                    val->SetInitializer(init.Get());
                 }
                 // Store the declaration so we can get the instruction to store too
                 scopes_.Set(v->name->symbol, val);
@@ -956,7 +949,7 @@
                 break;
         }
 
-        current_flow_block_->instructions.Push(inst);
+        current_flow_block_->Instructions().Push(inst);
         return inst;
     }
 
@@ -983,7 +976,7 @@
         BranchTo(if_node);
 
         auto* result = builder_.BlockParam(builder_.ir.types.Get<type::Bool>());
-        if_node->merge.target->As<Block>()->params.Push(result);
+        if_node->Merge().target->As<Block>()->SetParams(utils::Vector{result});
 
         utils::Result<Value*> rhs;
         {
@@ -997,17 +990,17 @@
             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->False().target->As<Block>();
+                BranchTo(if_node->Merge().target, std::move(alt_args));
 
-                current_flow_block_ = if_node->true_.target->As<Block>();
+                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->True().target->As<Block>();
+                BranchTo(if_node->Merge().target, std::move(alt_args));
 
-                current_flow_block_ = if_node->false_.target->As<Block>();
+                current_flow_block_ = if_node->False().target->As<Block>();
             }
 
             rhs = EmitExpression(expr->rhs);
@@ -1017,9 +1010,9 @@
             utils::Vector<Value*, 1> args;
             args.Push(rhs.Get());
 
-            BranchTo(if_node->merge.target, std::move(args));
+            BranchTo(if_node->Merge().target, std::move(args));
         }
-        current_flow_block_ = if_node->merge.target->As<Block>();
+        current_flow_block_ = if_node->Merge().target->As<Block>();
 
         return result;
     }
@@ -1101,7 +1094,7 @@
                 return utils::Failure;
         }
 
-        current_flow_block_->instructions.Push(inst);
+        current_flow_block_->Instructions().Push(inst);
         return inst;
     }
 
@@ -1115,7 +1108,7 @@
         auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx);
         auto* inst = builder_.Bitcast(ty, val.Get());
 
-        current_flow_block_->instructions.Push(inst);
+        current_flow_block_->Instructions().Push(inst);
         return inst;
     }
 
@@ -1178,7 +1171,7 @@
         if (inst == nullptr) {
             return utils::Failure;
         }
-        current_flow_block_->instructions.Push(inst);
+        current_flow_block_->Instructions().Push(inst);
         return inst;
     }
 
diff --git a/src/tint/ir/from_program_binary_test.cc b/src/tint/ir/from_program_binary_test.cc
index 85377ff..b74a8c8 100644
--- a/src/tint/ir/from_program_binary_test.cc
+++ b/src/tint/ir/from_program_binary_test.cc
@@ -830,7 +830,7 @@
     auto m = Build();
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool {
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:bool):bool {
   %fn2 = block {
   } -> %func_end true # return
 } %func_end
diff --git a/src/tint/ir/from_program_call_test.cc b/src/tint/ir/from_program_call_test.cc
index 038f4dd..6339c22 100644
--- a/src/tint/ir/from_program_call_test.cc
+++ b/src/tint/ir/from_program_call_test.cc
@@ -77,14 +77,14 @@
     auto m = Build();
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():void {
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:f32):void {
   %fn2 = block {
   } -> %func_end # return
 } %func_end
 
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
-    %1:void = call my_func, 6.0f
+    %2:void = call my_func, 6.0f
   } -> %func_end # return
 } %func_end
 
diff --git a/src/tint/ir/from_program_literal_test.cc b/src/tint/ir/from_program_literal_test.cc
index de5e81e..f746bc0 100644
--- a/src/tint/ir/from_program_literal_test.cc
+++ b/src/tint/ir/from_program_literal_test.cc
@@ -25,17 +25,19 @@
 namespace tint::ir {
 namespace {
 
-Value* GlobalVarInitializer(const Module& m) {
-    if (m.root_block->instructions.Length() == 0u) {
+const Value* GlobalVarInitializer(const Module& m) {
+    const auto instr = m.root_block->Instructions();
+
+    if (instr.Length() == 0u) {
         ADD_FAILURE() << "m.root_block has no instruction";
         return nullptr;
     }
-    auto* var = m.root_block->instructions[0]->As<ir::Var>();
+    auto* var = instr[0]->As<ir::Var>();
     if (!var) {
         ADD_FAILURE() << "m.root_block.instructions[0] was not a var";
         return nullptr;
     }
-    return var->initializer;
+    return var->Initializer();
 }
 
 using namespace tint::number_suffixes;  // NOLINT
@@ -51,7 +53,7 @@
 
     auto* init = GlobalVarInitializer(m.Get());
     ASSERT_TRUE(Is<Constant>(init));
-    auto* val = init->As<Constant>()->value;
+    auto* val = init->As<Constant>()->Value();
     EXPECT_TRUE(val->Is<constant::Scalar<bool>>());
     EXPECT_TRUE(val->As<constant::Scalar<bool>>()->ValueAs<bool>());
 }
@@ -65,11 +67,35 @@
 
     auto* init = GlobalVarInitializer(m.Get());
     ASSERT_TRUE(Is<Constant>(init));
-    auto* val = init->As<Constant>()->value;
+    auto* val = init->As<Constant>()->Value();
     EXPECT_TRUE(val->Is<constant::Scalar<bool>>());
     EXPECT_FALSE(val->As<constant::Scalar<bool>>()->ValueAs<bool>());
 }
 
+TEST_F(IR_BuilderImplTest, EmitLiteral_Bool_Deduped) {
+    GlobalVar("a", ty.bool_(), builtin::AddressSpace::kPrivate, Expr(true));
+    GlobalVar("b", ty.bool_(), builtin::AddressSpace::kPrivate, Expr(false));
+    GlobalVar("c", ty.bool_(), builtin::AddressSpace::kPrivate, Expr(true));
+    GlobalVar("d", ty.bool_(), builtin::AddressSpace::kPrivate, Expr(false));
+
+    auto m = Build();
+    ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+    auto instr = m.Get().root_block->Instructions();
+    auto* var_a = instr[0]->As<ir::Var>();
+    ASSERT_NE(var_a, nullptr);
+    auto* var_b = instr[1]->As<ir::Var>();
+    ASSERT_NE(var_b, nullptr);
+    auto* var_c = instr[2]->As<ir::Var>();
+    ASSERT_NE(var_c, nullptr);
+    auto* var_d = instr[3]->As<ir::Var>();
+    ASSERT_NE(var_d, nullptr);
+
+    ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
+    ASSERT_EQ(var_b->Initializer(), var_d->Initializer());
+    ASSERT_NE(var_a->Initializer(), var_b->Initializer());
+}
+
 TEST_F(IR_BuilderImplTest, EmitLiteral_F32) {
     auto* expr = Expr(1.2_f);
     GlobalVar("a", ty.f32(), builtin::AddressSpace::kPrivate, expr);
@@ -79,11 +105,31 @@
 
     auto* init = GlobalVarInitializer(m.Get());
     ASSERT_TRUE(Is<Constant>(init));
-    auto* val = init->As<Constant>()->value;
+    auto* val = init->As<Constant>()->Value();
     EXPECT_TRUE(val->Is<constant::Scalar<f32>>());
     EXPECT_EQ(1.2_f, val->As<constant::Scalar<f32>>()->ValueAs<f32>());
 }
 
+TEST_F(IR_BuilderImplTest, EmitLiteral_F32_Deduped) {
+    GlobalVar("a", ty.f32(), builtin::AddressSpace::kPrivate, Expr(1.2_f));
+    GlobalVar("b", ty.f32(), builtin::AddressSpace::kPrivate, Expr(1.25_f));
+    GlobalVar("c", ty.f32(), builtin::AddressSpace::kPrivate, Expr(1.2_f));
+
+    auto m = Build();
+    ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+    auto instr = m.Get().root_block->Instructions();
+    auto* var_a = instr[0]->As<ir::Var>();
+    ASSERT_NE(var_a, nullptr);
+    auto* var_b = instr[1]->As<ir::Var>();
+    ASSERT_NE(var_b, nullptr);
+    auto* var_c = instr[2]->As<ir::Var>();
+    ASSERT_NE(var_c, nullptr);
+
+    ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
+    ASSERT_NE(var_a->Initializer(), var_b->Initializer());
+}
+
 TEST_F(IR_BuilderImplTest, EmitLiteral_F16) {
     Enable(builtin::Extension::kF16);
     auto* expr = Expr(1.2_h);
@@ -94,11 +140,32 @@
 
     auto* init = GlobalVarInitializer(m.Get());
     ASSERT_TRUE(Is<Constant>(init));
-    auto* val = init->As<Constant>()->value;
+    auto* val = init->As<Constant>()->Value();
     EXPECT_TRUE(val->Is<constant::Scalar<f16>>());
     EXPECT_EQ(1.2_h, val->As<constant::Scalar<f16>>()->ValueAs<f32>());
 }
 
+TEST_F(IR_BuilderImplTest, EmitLiteral_F16_Deduped) {
+    Enable(builtin::Extension::kF16);
+    GlobalVar("a", ty.f16(), builtin::AddressSpace::kPrivate, Expr(1.2_h));
+    GlobalVar("b", ty.f16(), builtin::AddressSpace::kPrivate, Expr(1.25_h));
+    GlobalVar("c", ty.f16(), builtin::AddressSpace::kPrivate, Expr(1.2_h));
+
+    auto m = Build();
+    ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+    auto instr = m.Get().root_block->Instructions();
+    auto* var_a = instr[0]->As<ir::Var>();
+    ASSERT_NE(var_a, nullptr);
+    auto* var_b = instr[1]->As<ir::Var>();
+    ASSERT_NE(var_b, nullptr);
+    auto* var_c = instr[2]->As<ir::Var>();
+    ASSERT_NE(var_c, nullptr);
+
+    ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
+    ASSERT_NE(var_a->Initializer(), var_b->Initializer());
+}
+
 TEST_F(IR_BuilderImplTest, EmitLiteral_I32) {
     auto* expr = Expr(-2_i);
     GlobalVar("a", ty.i32(), builtin::AddressSpace::kPrivate, expr);
@@ -108,11 +175,31 @@
 
     auto* init = GlobalVarInitializer(m.Get());
     ASSERT_TRUE(Is<Constant>(init));
-    auto* val = init->As<Constant>()->value;
+    auto* val = init->As<Constant>()->Value();
     EXPECT_TRUE(val->Is<constant::Scalar<i32>>());
     EXPECT_EQ(-2_i, val->As<constant::Scalar<i32>>()->ValueAs<f32>());
 }
 
+TEST_F(IR_BuilderImplTest, EmitLiteral_I32_Deduped) {
+    GlobalVar("a", ty.i32(), builtin::AddressSpace::kPrivate, Expr(-2_i));
+    GlobalVar("b", ty.i32(), builtin::AddressSpace::kPrivate, Expr(2_i));
+    GlobalVar("c", ty.i32(), builtin::AddressSpace::kPrivate, Expr(-2_i));
+
+    auto m = Build();
+    ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+    auto instr = m.Get().root_block->Instructions();
+    auto* var_a = instr[0]->As<ir::Var>();
+    ASSERT_NE(var_a, nullptr);
+    auto* var_b = instr[1]->As<ir::Var>();
+    ASSERT_NE(var_b, nullptr);
+    auto* var_c = instr[2]->As<ir::Var>();
+    ASSERT_NE(var_c, nullptr);
+
+    ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
+    ASSERT_NE(var_a->Initializer(), var_b->Initializer());
+}
+
 TEST_F(IR_BuilderImplTest, EmitLiteral_U32) {
     auto* expr = Expr(2_u);
     GlobalVar("a", ty.u32(), builtin::AddressSpace::kPrivate, expr);
@@ -122,10 +209,30 @@
 
     auto* init = GlobalVarInitializer(m.Get());
     ASSERT_TRUE(Is<Constant>(init));
-    auto* val = init->As<Constant>()->value;
+    auto* val = init->As<Constant>()->Value();
     EXPECT_TRUE(val->Is<constant::Scalar<u32>>());
     EXPECT_EQ(2_u, val->As<constant::Scalar<u32>>()->ValueAs<f32>());
 }
 
+TEST_F(IR_BuilderImplTest, EmitLiteral_U32_Deduped) {
+    GlobalVar("a", ty.u32(), builtin::AddressSpace::kPrivate, Expr(2_u));
+    GlobalVar("b", ty.u32(), builtin::AddressSpace::kPrivate, Expr(3_u));
+    GlobalVar("c", ty.u32(), builtin::AddressSpace::kPrivate, Expr(2_u));
+
+    auto m = Build();
+    ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+    auto instr = m.Get().root_block->Instructions();
+    auto* var_a = instr[0]->As<ir::Var>();
+    ASSERT_NE(var_a, nullptr);
+    auto* var_b = instr[1]->As<ir::Var>();
+    ASSERT_NE(var_b, nullptr);
+    auto* var_c = instr[2]->As<ir::Var>();
+    ASSERT_NE(var_c, nullptr);
+
+    ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
+    ASSERT_NE(var_a->Initializer(), var_b->Initializer());
+}
+
 }  // namespace
 }  // namespace tint::ir
diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc
index 47dc911..1049386 100644
--- a/src/tint/ir/from_program_test.cc
+++ b/src/tint/ir/from_program_test.cc
@@ -63,13 +63,13 @@
     ASSERT_EQ(1u, m->functions.Length());
 
     auto* f = m->functions[0];
-    ASSERT_NE(f->start_target, nullptr);
-    ASSERT_NE(f->end_target, nullptr);
+    ASSERT_NE(f->StartTarget(), nullptr);
+    ASSERT_NE(f->EndTarget(), nullptr);
 
-    EXPECT_EQ(1u, f->start_target->inbound_branches.Length());
-    EXPECT_EQ(1u, f->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, f->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length());
 
-    EXPECT_EQ(m->functions[0]->pipeline_stage, Function::PipelineStage::kUndefined);
+    EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined);
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f():void {
   %fn2 = block {
@@ -79,6 +79,57 @@
 )");
 }
 
+TEST_F(IR_BuilderImplTest, Func_WithParam) {
+    Func("f", utils::Vector{Param("a", ty.u32())}, ty.u32(), utils::Vector{Return("a")});
+
+    auto m = Build();
+    ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+    ASSERT_EQ(1u, m->functions.Length());
+
+    auto* f = m->functions[0];
+    ASSERT_NE(f->StartTarget(), nullptr);
+    ASSERT_NE(f->EndTarget(), nullptr);
+
+    EXPECT_EQ(1u, f->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length());
+
+    EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined);
+
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f(%a:u32):u32 {
+  %fn2 = block {
+  } -> %func_end %a # return
+} %func_end
+
+)");
+}
+
+TEST_F(IR_BuilderImplTest, Func_WithMultipleParam) {
+    Func("f", utils::Vector{Param("a", ty.u32()), Param("b", ty.i32()), Param("c", ty.bool_())},
+         ty.void_(), utils::Empty);
+
+    auto m = Build();
+    ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+    ASSERT_EQ(1u, m->functions.Length());
+
+    auto* f = m->functions[0];
+    ASSERT_NE(f->StartTarget(), nullptr);
+    ASSERT_NE(f->EndTarget(), nullptr);
+
+    EXPECT_EQ(1u, f->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length());
+
+    EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined);
+
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f(%a:u32, %b:i32, %c:bool):void {
+  %fn2 = block {
+  } -> %func_end # return
+} %func_end
+
+)");
+}
+
 TEST_F(IR_BuilderImplTest, EntryPoint) {
     Func("f", utils::Empty, ty.void_(), utils::Empty,
          utils::Vector{Stage(ast::PipelineStage::kFragment)});
@@ -86,7 +137,7 @@
     auto m = Build();
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
-    EXPECT_EQ(m->functions[0]->pipeline_stage, Function::PipelineStage::kFragment);
+    EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kFragment);
 }
 
 TEST_F(IR_BuilderImplTest, IfStatement) {
@@ -97,19 +148,19 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::If>(m.Get());
-    ASSERT_NE(flow->true_.target, nullptr);
-    ASSERT_NE(flow->false_.target, nullptr);
-    ASSERT_NE(flow->merge.target, nullptr);
+    ASSERT_NE(flow->True().target, nullptr);
+    ASSERT_NE(flow->False().target, nullptr);
+    ASSERT_NE(flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->false_.target->inbound_branches.Length());
-    EXPECT_EQ(2u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->False().target->InboundBranches().Length());
+    EXPECT_EQ(2u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -141,19 +192,19 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::If>(m.Get());
-    ASSERT_NE(flow->true_.target, nullptr);
-    ASSERT_NE(flow->false_.target, nullptr);
-    ASSERT_NE(flow->merge.target, nullptr);
+    ASSERT_NE(flow->True().target, nullptr);
+    ASSERT_NE(flow->False().target, nullptr);
+    ASSERT_NE(flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->false_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
-    EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->False().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -184,19 +235,19 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::If>(m.Get());
-    ASSERT_NE(flow->true_.target, nullptr);
-    ASSERT_NE(flow->false_.target, nullptr);
-    ASSERT_NE(flow->merge.target, nullptr);
+    ASSERT_NE(flow->True().target, nullptr);
+    ASSERT_NE(flow->False().target, nullptr);
+    ASSERT_NE(flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->false_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
-    EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->False().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -227,19 +278,19 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::If>(m.Get());
-    ASSERT_NE(flow->true_.target, nullptr);
-    ASSERT_NE(flow->false_.target, nullptr);
-    ASSERT_NE(flow->merge.target, nullptr);
+    ASSERT_NE(flow->True().target, nullptr);
+    ASSERT_NE(flow->False().target, nullptr);
+    ASSERT_NE(flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->false_.target->inbound_branches.Length());
-    EXPECT_EQ(0u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
-    EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->False().target->InboundBranches().Length());
+    EXPECT_EQ(0u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -267,15 +318,15 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* if_flow = FindSingleFlowNode<ir::If>(m.Get());
-    ASSERT_NE(if_flow->true_.target, nullptr);
-    ASSERT_NE(if_flow->false_.target, nullptr);
-    ASSERT_NE(if_flow->merge.target, nullptr);
+    ASSERT_NE(if_flow->True().target, nullptr);
+    ASSERT_NE(if_flow->False().target, nullptr);
+    ASSERT_NE(if_flow->Merge().target, nullptr);
 
     auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
     ASSERT_NE(loop_flow, nullptr);
-    ASSERT_NE(loop_flow->start.target, nullptr);
-    ASSERT_NE(loop_flow->continuing.target, nullptr);
-    ASSERT_NE(loop_flow->merge.target, nullptr);
+    ASSERT_NE(loop_flow->Start().target, nullptr);
+    ASSERT_NE(loop_flow->Continuing().target, nullptr);
+    ASSERT_NE(loop_flow->Merge().target, nullptr);
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -316,19 +367,19 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::Loop>(m.Get());
-    ASSERT_NE(flow->start.target, nullptr);
-    ASSERT_NE(flow->continuing.target, nullptr);
-    ASSERT_NE(flow->merge.target, nullptr);
+    ASSERT_NE(flow->Start().target, nullptr);
+    ASSERT_NE(flow->Continuing().target, nullptr);
+    ASSERT_NE(flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(2u, flow->start.target->inbound_branches.Length());
-    EXPECT_EQ(0u, flow->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length());
+    EXPECT_EQ(0u, flow->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -357,28 +408,28 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
-    ASSERT_NE(loop_flow->start.target, nullptr);
-    ASSERT_NE(loop_flow->continuing.target, nullptr);
-    ASSERT_NE(loop_flow->merge.target, nullptr);
+    ASSERT_NE(loop_flow->Start().target, nullptr);
+    ASSERT_NE(loop_flow->Continuing().target, nullptr);
+    ASSERT_NE(loop_flow->Merge().target, nullptr);
 
     auto* if_flow = FindSingleFlowNode<ir::If>(m.Get());
-    ASSERT_NE(if_flow->true_.target, nullptr);
-    ASSERT_NE(if_flow->false_.target, nullptr);
-    ASSERT_NE(if_flow->merge.target, nullptr);
+    ASSERT_NE(if_flow->True().target, nullptr);
+    ASSERT_NE(if_flow->False().target, nullptr);
+    ASSERT_NE(if_flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, loop_flow->inbound_branches.Length());
-    EXPECT_EQ(2u, loop_flow->start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, loop_flow->InboundBranches().Length());
+    EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -424,28 +475,28 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
-    ASSERT_NE(loop_flow->start.target, nullptr);
-    ASSERT_NE(loop_flow->continuing.target, nullptr);
-    ASSERT_NE(loop_flow->merge.target, nullptr);
+    ASSERT_NE(loop_flow->Start().target, nullptr);
+    ASSERT_NE(loop_flow->Continuing().target, nullptr);
+    ASSERT_NE(loop_flow->Merge().target, nullptr);
 
     auto* break_if_flow = FindSingleFlowNode<ir::If>(m.Get());
-    ASSERT_NE(break_if_flow->true_.target, nullptr);
-    ASSERT_NE(break_if_flow->false_.target, nullptr);
-    ASSERT_NE(break_if_flow->merge.target, nullptr);
+    ASSERT_NE(break_if_flow->True().target, nullptr);
+    ASSERT_NE(break_if_flow->False().target, nullptr);
+    ASSERT_NE(break_if_flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, loop_flow->inbound_branches.Length());
-    EXPECT_EQ(2u, loop_flow->start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, break_if_flow->inbound_branches.Length());
-    EXPECT_EQ(1u, break_if_flow->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, break_if_flow->false_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, break_if_flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, loop_flow->InboundBranches().Length());
+    EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, break_if_flow->InboundBranches().Length());
+    EXPECT_EQ(1u, break_if_flow->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, break_if_flow->False().target->InboundBranches().Length());
+    EXPECT_EQ(1u, break_if_flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -535,28 +586,28 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
-    ASSERT_NE(loop_flow->start.target, nullptr);
-    ASSERT_NE(loop_flow->continuing.target, nullptr);
-    ASSERT_NE(loop_flow->merge.target, nullptr);
+    ASSERT_NE(loop_flow->Start().target, nullptr);
+    ASSERT_NE(loop_flow->Continuing().target, nullptr);
+    ASSERT_NE(loop_flow->Merge().target, nullptr);
 
     auto* if_flow = FindSingleFlowNode<ir::If>(m.Get());
-    ASSERT_NE(if_flow->true_.target, nullptr);
-    ASSERT_NE(if_flow->false_.target, nullptr);
-    ASSERT_NE(if_flow->merge.target, nullptr);
+    ASSERT_NE(if_flow->True().target, nullptr);
+    ASSERT_NE(if_flow->False().target, nullptr);
+    ASSERT_NE(if_flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, loop_flow->inbound_branches.Length());
-    EXPECT_EQ(2u, loop_flow->start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(0u, loop_flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, loop_flow->InboundBranches().Length());
+    EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(0u, loop_flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -597,19 +648,19 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
-    ASSERT_NE(loop_flow->start.target, nullptr);
-    ASSERT_NE(loop_flow->continuing.target, nullptr);
-    ASSERT_NE(loop_flow->merge.target, nullptr);
+    ASSERT_NE(loop_flow->Start().target, nullptr);
+    ASSERT_NE(loop_flow->Continuing().target, nullptr);
+    ASSERT_NE(loop_flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, loop_flow->inbound_branches.Length());
-    EXPECT_EQ(2u, loop_flow->start.target->inbound_branches.Length());
-    EXPECT_EQ(0u, loop_flow->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(0u, loop_flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, loop_flow->InboundBranches().Length());
+    EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length());
+    EXPECT_EQ(0u, loop_flow->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(0u, loop_flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -642,25 +693,25 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
-    ASSERT_NE(loop_flow->start.target, nullptr);
-    ASSERT_NE(loop_flow->continuing.target, nullptr);
-    ASSERT_NE(loop_flow->merge.target, nullptr);
+    ASSERT_NE(loop_flow->Start().target, nullptr);
+    ASSERT_NE(loop_flow->Continuing().target, nullptr);
+    ASSERT_NE(loop_flow->Merge().target, nullptr);
 
     auto* break_if_flow = FindSingleFlowNode<ir::If>(m.Get());
-    ASSERT_NE(break_if_flow->true_.target, nullptr);
-    ASSERT_NE(break_if_flow->false_.target, nullptr);
-    ASSERT_NE(break_if_flow->merge.target, nullptr);
+    ASSERT_NE(break_if_flow->True().target, nullptr);
+    ASSERT_NE(break_if_flow->False().target, nullptr);
+    ASSERT_NE(break_if_flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, loop_flow->inbound_branches.Length());
-    EXPECT_EQ(2u, loop_flow->start.target->inbound_branches.Length());
-    EXPECT_EQ(0u, loop_flow->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
+    EXPECT_EQ(1u, loop_flow->InboundBranches().Length());
+    EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length());
+    EXPECT_EQ(0u, loop_flow->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
     // This is 1 because only the loop branch happens. The subsequent if return is dead code.
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -685,28 +736,28 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
-    ASSERT_NE(loop_flow->start.target, nullptr);
-    ASSERT_NE(loop_flow->continuing.target, nullptr);
-    ASSERT_NE(loop_flow->merge.target, nullptr);
+    ASSERT_NE(loop_flow->Start().target, nullptr);
+    ASSERT_NE(loop_flow->Continuing().target, nullptr);
+    ASSERT_NE(loop_flow->Merge().target, nullptr);
 
     auto* if_flow = FindSingleFlowNode<ir::If>(m.Get());
-    ASSERT_NE(if_flow->true_.target, nullptr);
-    ASSERT_NE(if_flow->false_.target, nullptr);
-    ASSERT_NE(if_flow->merge.target, nullptr);
+    ASSERT_NE(if_flow->True().target, nullptr);
+    ASSERT_NE(if_flow->False().target, nullptr);
+    ASSERT_NE(if_flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, loop_flow->inbound_branches.Length());
-    EXPECT_EQ(2u, loop_flow->start.target->inbound_branches.Length());
-    EXPECT_EQ(0u, loop_flow->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(2u, loop_flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
-    EXPECT_EQ(0u, if_flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, loop_flow->InboundBranches().Length());
+    EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length());
+    EXPECT_EQ(0u, loop_flow->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(2u, loop_flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length());
+    EXPECT_EQ(0u, if_flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -756,96 +807,96 @@
 
     auto block_exit = [&](const ir::FlowNode* node) -> const ir::FlowNode* {
         if (auto* block = As<ir::Block>(node)) {
-            return block->branch.target;
+            return block->Branch().target;
         }
         return nullptr;
     };
 
-    auto* loop_flow_a = As<ir::Loop>(m->functions[0]->start_target->branch.target);
+    auto* loop_flow_a = As<ir::Loop>(m->functions[0]->StartTarget()->Branch().target);
     ASSERT_NE(loop_flow_a, nullptr);
-    ASSERT_NE(loop_flow_a->start.target, nullptr);
-    ASSERT_NE(loop_flow_a->continuing.target, nullptr);
-    ASSERT_NE(loop_flow_a->merge.target, nullptr);
+    ASSERT_NE(loop_flow_a->Start().target, nullptr);
+    ASSERT_NE(loop_flow_a->Continuing().target, nullptr);
+    ASSERT_NE(loop_flow_a->Merge().target, nullptr);
 
-    auto* loop_flow_b = As<ir::Loop>(block_exit(loop_flow_a->start.target));
+    auto* loop_flow_b = As<ir::Loop>(block_exit(loop_flow_a->Start().target));
     ASSERT_NE(loop_flow_b, nullptr);
-    ASSERT_NE(loop_flow_b->start.target, nullptr);
-    ASSERT_NE(loop_flow_b->continuing.target, nullptr);
-    ASSERT_NE(loop_flow_b->merge.target, nullptr);
+    ASSERT_NE(loop_flow_b->Start().target, nullptr);
+    ASSERT_NE(loop_flow_b->Continuing().target, nullptr);
+    ASSERT_NE(loop_flow_b->Merge().target, nullptr);
 
-    auto* if_flow_a = As<ir::If>(block_exit(loop_flow_b->start.target));
+    auto* if_flow_a = As<ir::If>(block_exit(loop_flow_b->Start().target));
     ASSERT_NE(if_flow_a, nullptr);
-    ASSERT_NE(if_flow_a->true_.target, nullptr);
-    ASSERT_NE(if_flow_a->false_.target, nullptr);
-    ASSERT_NE(if_flow_a->merge.target, nullptr);
+    ASSERT_NE(if_flow_a->True().target, nullptr);
+    ASSERT_NE(if_flow_a->False().target, nullptr);
+    ASSERT_NE(if_flow_a->Merge().target, nullptr);
 
-    auto* if_flow_b = As<ir::If>(block_exit(if_flow_a->merge.target));
+    auto* if_flow_b = As<ir::If>(block_exit(if_flow_a->Merge().target));
     ASSERT_NE(if_flow_b, nullptr);
-    ASSERT_NE(if_flow_b->true_.target, nullptr);
-    ASSERT_NE(if_flow_b->false_.target, nullptr);
-    ASSERT_NE(if_flow_b->merge.target, nullptr);
+    ASSERT_NE(if_flow_b->True().target, nullptr);
+    ASSERT_NE(if_flow_b->False().target, nullptr);
+    ASSERT_NE(if_flow_b->Merge().target, nullptr);
 
-    auto* loop_flow_c = As<ir::Loop>(block_exit(loop_flow_b->continuing.target));
+    auto* loop_flow_c = As<ir::Loop>(block_exit(loop_flow_b->Continuing().target));
     ASSERT_NE(loop_flow_c, nullptr);
-    ASSERT_NE(loop_flow_c->start.target, nullptr);
-    ASSERT_NE(loop_flow_c->continuing.target, nullptr);
-    ASSERT_NE(loop_flow_c->merge.target, nullptr);
+    ASSERT_NE(loop_flow_c->Start().target, nullptr);
+    ASSERT_NE(loop_flow_c->Continuing().target, nullptr);
+    ASSERT_NE(loop_flow_c->Merge().target, nullptr);
 
-    auto* loop_flow_d = As<ir::Loop>(block_exit(loop_flow_c->merge.target));
+    auto* loop_flow_d = As<ir::Loop>(block_exit(loop_flow_c->Merge().target));
     ASSERT_NE(loop_flow_d, nullptr);
-    ASSERT_NE(loop_flow_d->start.target, nullptr);
-    ASSERT_NE(loop_flow_d->continuing.target, nullptr);
-    ASSERT_NE(loop_flow_d->merge.target, nullptr);
+    ASSERT_NE(loop_flow_d->Start().target, nullptr);
+    ASSERT_NE(loop_flow_d->Continuing().target, nullptr);
+    ASSERT_NE(loop_flow_d->Merge().target, nullptr);
 
-    auto* if_flow_c = As<ir::If>(block_exit(loop_flow_d->continuing.target));
+    auto* if_flow_c = As<ir::If>(block_exit(loop_flow_d->Continuing().target));
     ASSERT_NE(if_flow_c, nullptr);
-    ASSERT_NE(if_flow_c->true_.target, nullptr);
-    ASSERT_NE(if_flow_c->false_.target, nullptr);
-    ASSERT_NE(if_flow_c->merge.target, nullptr);
+    ASSERT_NE(if_flow_c->True().target, nullptr);
+    ASSERT_NE(if_flow_c->False().target, nullptr);
+    ASSERT_NE(if_flow_c->Merge().target, nullptr);
 
-    auto* if_flow_d = As<ir::If>(block_exit(loop_flow_b->merge.target));
+    auto* if_flow_d = As<ir::If>(block_exit(loop_flow_b->Merge().target));
     ASSERT_NE(if_flow_d, nullptr);
-    ASSERT_NE(if_flow_d->true_.target, nullptr);
-    ASSERT_NE(if_flow_d->false_.target, nullptr);
-    ASSERT_NE(if_flow_d->merge.target, nullptr);
+    ASSERT_NE(if_flow_d->True().target, nullptr);
+    ASSERT_NE(if_flow_d->False().target, nullptr);
+    ASSERT_NE(if_flow_d->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, loop_flow_a->inbound_branches.Length());
-    EXPECT_EQ(2u, loop_flow_a->start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow_a->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow_a->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow_b->inbound_branches.Length());
-    EXPECT_EQ(2u, loop_flow_b->start.target->inbound_branches.Length());
-    EXPECT_EQ(2u, loop_flow_b->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow_b->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow_c->inbound_branches.Length());
-    EXPECT_EQ(2u, loop_flow_c->start.target->inbound_branches.Length());
-    EXPECT_EQ(0u, loop_flow_c->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow_c->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow_d->inbound_branches.Length());
-    EXPECT_EQ(2u, loop_flow_d->start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow_d->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(1u, loop_flow_d->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_a->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_a->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_a->false_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_a->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_b->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_b->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_b->false_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_b->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_c->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_c->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_c->false_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_c->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_d->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_d->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_d->false_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow_d->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, loop_flow_a->InboundBranches().Length());
+    EXPECT_EQ(2u, loop_flow_a->Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow_a->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow_a->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow_b->InboundBranches().Length());
+    EXPECT_EQ(2u, loop_flow_b->Start().target->InboundBranches().Length());
+    EXPECT_EQ(2u, loop_flow_b->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow_b->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow_c->InboundBranches().Length());
+    EXPECT_EQ(2u, loop_flow_c->Start().target->InboundBranches().Length());
+    EXPECT_EQ(0u, loop_flow_c->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow_c->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow_d->InboundBranches().Length());
+    EXPECT_EQ(2u, loop_flow_d->Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow_d->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow_d->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_a->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_a->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_a->False().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_a->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_b->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_b->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_b->False().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_b->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_c->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_c->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_c->False().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_c->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_d->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_d->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_d->False().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow_d->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -964,28 +1015,28 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::Loop>(m.Get());
-    ASSERT_NE(flow->start.target, nullptr);
-    ASSERT_NE(flow->continuing.target, nullptr);
-    ASSERT_NE(flow->merge.target, nullptr);
+    ASSERT_NE(flow->Start().target, nullptr);
+    ASSERT_NE(flow->Continuing().target, nullptr);
+    ASSERT_NE(flow->Merge().target, nullptr);
 
-    ASSERT_NE(flow->start.target->As<ir::Block>()->branch.target, nullptr);
-    ASSERT_TRUE(flow->start.target->As<ir::Block>()->branch.target->Is<ir::If>());
-    auto* if_flow = flow->start.target->As<ir::Block>()->branch.target->As<ir::If>();
-    ASSERT_NE(if_flow->true_.target, nullptr);
-    ASSERT_NE(if_flow->false_.target, nullptr);
-    ASSERT_NE(if_flow->merge.target, nullptr);
+    ASSERT_NE(flow->Start().target->As<ir::Block>()->Branch().target, nullptr);
+    ASSERT_TRUE(flow->Start().target->As<ir::Block>()->Branch().target->Is<ir::If>());
+    auto* if_flow = flow->Start().target->As<ir::Block>()->Branch().target->As<ir::If>();
+    ASSERT_NE(if_flow->True().target, nullptr);
+    ASSERT_NE(if_flow->False().target, nullptr);
+    ASSERT_NE(if_flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(2u, flow->start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1030,28 +1081,28 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::Loop>(m.Get());
-    ASSERT_NE(flow->start.target, nullptr);
-    ASSERT_NE(flow->continuing.target, nullptr);
-    ASSERT_NE(flow->merge.target, nullptr);
+    ASSERT_NE(flow->Start().target, nullptr);
+    ASSERT_NE(flow->Continuing().target, nullptr);
+    ASSERT_NE(flow->Merge().target, nullptr);
 
-    ASSERT_NE(flow->start.target->As<ir::Block>()->branch.target, nullptr);
-    ASSERT_TRUE(flow->start.target->As<ir::Block>()->branch.target->Is<ir::If>());
-    auto* if_flow = flow->start.target->As<ir::Block>()->branch.target->As<ir::If>();
-    ASSERT_NE(if_flow->true_.target, nullptr);
-    ASSERT_NE(if_flow->false_.target, nullptr);
-    ASSERT_NE(if_flow->merge.target, nullptr);
+    ASSERT_NE(flow->Start().target->As<ir::Block>()->Branch().target, nullptr);
+    ASSERT_TRUE(flow->Start().target->As<ir::Block>()->Branch().target->Is<ir::If>());
+    auto* if_flow = flow->Start().target->As<ir::Block>()->Branch().target->As<ir::If>();
+    ASSERT_NE(if_flow->True().target, nullptr);
+    ASSERT_NE(if_flow->False().target, nullptr);
+    ASSERT_NE(if_flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(2u, flow->start.target->inbound_branches.Length());
-    EXPECT_EQ(0u, flow->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
+    EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length());
+    EXPECT_EQ(0u, flow->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1104,28 +1155,28 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::Loop>(m.Get());
-    ASSERT_NE(flow->start.target, nullptr);
-    ASSERT_NE(flow->continuing.target, nullptr);
-    ASSERT_NE(flow->merge.target, nullptr);
+    ASSERT_NE(flow->Start().target, nullptr);
+    ASSERT_NE(flow->Continuing().target, nullptr);
+    ASSERT_NE(flow->Merge().target, nullptr);
 
-    ASSERT_NE(flow->start.target->As<ir::Block>()->branch.target, nullptr);
-    ASSERT_TRUE(flow->start.target->As<ir::Block>()->branch.target->Is<ir::If>());
-    auto* if_flow = flow->start.target->As<ir::Block>()->branch.target->As<ir::If>();
-    ASSERT_NE(if_flow->true_.target, nullptr);
-    ASSERT_NE(if_flow->false_.target, nullptr);
-    ASSERT_NE(if_flow->merge.target, nullptr);
+    ASSERT_NE(flow->Start().target->As<ir::Block>()->Branch().target, nullptr);
+    ASSERT_TRUE(flow->Start().target->As<ir::Block>()->Branch().target->Is<ir::If>());
+    auto* if_flow = flow->Start().target->As<ir::Block>()->Branch().target->As<ir::If>();
+    ASSERT_NE(if_flow->True().target, nullptr);
+    ASSERT_NE(if_flow->False().target, nullptr);
+    ASSERT_NE(if_flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(2u, flow->start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->true_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
-    EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()), R"()");
 }
@@ -1138,18 +1189,18 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::Loop>(m.Get());
-    ASSERT_NE(flow->start.target, nullptr);
-    ASSERT_NE(flow->continuing.target, nullptr);
-    ASSERT_NE(flow->merge.target, nullptr);
+    ASSERT_NE(flow->Start().target, nullptr);
+    ASSERT_NE(flow->Continuing().target, nullptr);
+    ASSERT_NE(flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(2u, flow->start.target->inbound_branches.Length());
-    EXPECT_EQ(0u, flow->continuing.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length());
+    EXPECT_EQ(0u, flow->Continuing().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1180,31 +1231,33 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::Switch>(m.Get());
-    ASSERT_NE(flow->merge.target, nullptr);
-    ASSERT_EQ(3u, flow->cases.Length());
+    ASSERT_NE(flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    ASSERT_EQ(1u, flow->cases[0].selectors.Length());
-    ASSERT_TRUE(flow->cases[0].selectors[0].val->value->Is<constant::Scalar<tint::i32>>());
+    auto cases = flow->Cases();
+    ASSERT_EQ(3u, cases.Length());
+
+    ASSERT_EQ(1u, cases[0].selectors.Length());
+    ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
     EXPECT_EQ(0_i,
-              flow->cases[0].selectors[0].val->value->As<constant::Scalar<tint::i32>>()->ValueOf());
+              cases[0].selectors[0].val->Value()->As<constant::Scalar<tint::i32>>()->ValueOf());
 
-    ASSERT_EQ(1u, flow->cases[1].selectors.Length());
-    ASSERT_TRUE(flow->cases[1].selectors[0].val->value->Is<constant::Scalar<tint::i32>>());
+    ASSERT_EQ(1u, cases[1].selectors.Length());
+    ASSERT_TRUE(cases[1].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
     EXPECT_EQ(1_i,
-              flow->cases[1].selectors[0].val->value->As<constant::Scalar<tint::i32>>()->ValueOf());
+              cases[1].selectors[0].val->Value()->As<constant::Scalar<tint::i32>>()->ValueOf());
 
-    ASSERT_EQ(1u, flow->cases[2].selectors.Length());
-    EXPECT_TRUE(flow->cases[2].selectors[0].IsDefault());
+    ASSERT_EQ(1u, cases[2].selectors.Length());
+    EXPECT_TRUE(cases[2].selectors[0].IsDefault());
 
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->cases[0].start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->cases[1].start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->cases[2].start.target->inbound_branches.Length());
-    EXPECT_EQ(3u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, cases[1].Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, cases[2].Start().target->InboundBranches().Length());
+    EXPECT_EQ(3u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1244,27 +1297,28 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::Switch>(m.Get());
-    ASSERT_NE(flow->merge.target, nullptr);
-    ASSERT_EQ(1u, flow->cases.Length());
+    ASSERT_NE(flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    ASSERT_EQ(3u, flow->cases[0].selectors.Length());
-    ASSERT_TRUE(flow->cases[0].selectors[0].val->value->Is<constant::Scalar<tint::i32>>());
+    auto cases = flow->Cases();
+    ASSERT_EQ(1u, cases.Length());
+    ASSERT_EQ(3u, cases[0].selectors.Length());
+    ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
     EXPECT_EQ(0_i,
-              flow->cases[0].selectors[0].val->value->As<constant::Scalar<tint::i32>>()->ValueOf());
+              cases[0].selectors[0].val->Value()->As<constant::Scalar<tint::i32>>()->ValueOf());
 
-    ASSERT_TRUE(flow->cases[0].selectors[1].val->value->Is<constant::Scalar<tint::i32>>());
+    ASSERT_TRUE(cases[0].selectors[1].val->Value()->Is<constant::Scalar<tint::i32>>());
     EXPECT_EQ(1_i,
-              flow->cases[0].selectors[1].val->value->As<constant::Scalar<tint::i32>>()->ValueOf());
+              cases[0].selectors[1].val->Value()->As<constant::Scalar<tint::i32>>()->ValueOf());
 
-    EXPECT_TRUE(flow->cases[0].selectors[2].IsDefault());
+    EXPECT_TRUE(cases[0].selectors[2].IsDefault());
 
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->cases[0].start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1292,19 +1346,20 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::Switch>(m.Get());
-    ASSERT_NE(flow->merge.target, nullptr);
-    ASSERT_EQ(1u, flow->cases.Length());
+    ASSERT_NE(flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    ASSERT_EQ(1u, flow->cases[0].selectors.Length());
-    EXPECT_TRUE(flow->cases[0].selectors[0].IsDefault());
+    auto cases = flow->Cases();
+    ASSERT_EQ(1u, cases.Length());
+    ASSERT_EQ(1u, cases[0].selectors.Length());
+    EXPECT_TRUE(cases[0].selectors[0].IsDefault());
 
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->cases[0].start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1334,26 +1389,27 @@
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     auto* flow = FindSingleFlowNode<ir::Switch>(m.Get());
-    ASSERT_NE(flow->merge.target, nullptr);
-    ASSERT_EQ(2u, flow->cases.Length());
+    ASSERT_NE(flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    ASSERT_EQ(1u, flow->cases[0].selectors.Length());
-    ASSERT_TRUE(flow->cases[0].selectors[0].val->value->Is<constant::Scalar<tint::i32>>());
+    auto cases = flow->Cases();
+    ASSERT_EQ(2u, cases.Length());
+    ASSERT_EQ(1u, cases[0].selectors.Length());
+    ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
     EXPECT_EQ(0_i,
-              flow->cases[0].selectors[0].val->value->As<constant::Scalar<tint::i32>>()->ValueOf());
+              cases[0].selectors[0].val->Value()->As<constant::Scalar<tint::i32>>()->ValueOf());
 
-    ASSERT_EQ(1u, flow->cases[1].selectors.Length());
-    EXPECT_TRUE(flow->cases[1].selectors[0].IsDefault());
+    ASSERT_EQ(1u, cases[1].selectors.Length());
+    EXPECT_TRUE(cases[1].selectors[0].IsDefault());
 
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->cases[0].start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->cases[1].start.target->inbound_branches.Length());
-    EXPECT_EQ(2u, flow->merge.target->inbound_branches.Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, cases[1].Start().target->InboundBranches().Length());
+    EXPECT_EQ(2u, flow->Merge().target->InboundBranches().Length());
     // This is 1 because the if is dead-code eliminated and the return doesn't happen.
-    EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1390,25 +1446,26 @@
     ASSERT_EQ(FindSingleFlowNode<ir::If>(m.Get()), nullptr);
 
     auto* flow = FindSingleFlowNode<ir::Switch>(m.Get());
-    ASSERT_NE(flow->merge.target, nullptr);
-    ASSERT_EQ(2u, flow->cases.Length());
+    ASSERT_NE(flow->Merge().target, nullptr);
 
     ASSERT_EQ(1u, m->functions.Length());
     auto* func = m->functions[0];
 
-    ASSERT_EQ(1u, flow->cases[0].selectors.Length());
-    ASSERT_TRUE(flow->cases[0].selectors[0].val->value->Is<constant::Scalar<tint::i32>>());
+    auto cases = flow->Cases();
+    ASSERT_EQ(2u, cases.Length());
+    ASSERT_EQ(1u, cases[0].selectors.Length());
+    ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
     EXPECT_EQ(0_i,
-              flow->cases[0].selectors[0].val->value->As<constant::Scalar<tint::i32>>()->ValueOf());
+              cases[0].selectors[0].val->Value()->As<constant::Scalar<tint::i32>>()->ValueOf());
 
-    ASSERT_EQ(1u, flow->cases[1].selectors.Length());
-    EXPECT_TRUE(flow->cases[1].selectors[0].IsDefault());
+    ASSERT_EQ(1u, cases[1].selectors.Length());
+    EXPECT_TRUE(cases[1].selectors[0].IsDefault());
 
-    EXPECT_EQ(1u, flow->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->cases[0].start.target->inbound_branches.Length());
-    EXPECT_EQ(1u, flow->cases[1].start.target->inbound_branches.Length());
-    EXPECT_EQ(0u, flow->merge.target->inbound_branches.Length());
-    EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
+    EXPECT_EQ(1u, flow->InboundBranches().Length());
+    EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length());
+    EXPECT_EQ(1u, cases[1].Start().target->InboundBranches().Length());
+    EXPECT_EQ(0u, flow->Merge().target->InboundBranches().Length());
+    EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
diff --git a/src/tint/ir/function.cc b/src/tint/ir/function.cc
index 2a5f8c3..d6cd6fc 100644
--- a/src/tint/ir/function.cc
+++ b/src/tint/ir/function.cc
@@ -18,11 +18,11 @@
 
 namespace tint::ir {
 
-Function::Function(Symbol n,
+Function::Function(Symbol name,
                    type::Type* rt,
                    PipelineStage stage,
                    std::optional<std::array<uint32_t, 3>> wg_size)
-    : Base(), name(n), pipeline_stage(stage), workgroup_size(wg_size), return_type(rt) {}
+    : Base(), name_(name), return_type_(rt), pipeline_stage_(stage), workgroup_size_(wg_size) {}
 
 Function::~Function() = default;
 
diff --git a/src/tint/ir/function.h b/src/tint/ir/function.h
index 95487b9..86465aa 100644
--- a/src/tint/ir/function.h
+++ b/src/tint/ir/function.h
@@ -17,8 +17,10 @@
 
 #include <array>
 #include <optional>
+#include <utility>
 
 #include "src/tint/ir/flow_node.h"
+#include "src/tint/ir/function_param.h"
 #include "src/tint/symbol.h"
 #include "src/tint/type/type.h"
 
@@ -70,29 +72,81 @@
              type::Type* rt,
              PipelineStage stage = PipelineStage::kUndefined,
              std::optional<std::array<uint32_t, 3>> wg_size = {});
+    Function(Function&&) = delete;
+    Function(const Function&) = delete;
     ~Function() override;
 
-    /// The function name
-    Symbol name;
+    Function& operator=(Function&&) = delete;
+    Function& operator=(const Function&) = delete;
 
-    /// The pipeline stage for the function, `kUndefined` if the function is not an entry point
-    PipelineStage pipeline_stage = PipelineStage::kUndefined;
+    /// @returns the function name
+    Symbol Name() const { return name_; }
 
-    /// If this is a `compute` entry point, holds the workgroup size information
-    std::optional<std::array<uint32_t, 3>> workgroup_size;
+    /// Sets the function stage
+    /// @param stage the stage to set
+    void SetStage(PipelineStage stage) { pipeline_stage_ = stage; }
 
-    /// The function return type
-    const type::Type* return_type = nullptr;
-    /// The function return attributes if any
-    utils::Vector<ReturnAttribute, 1> return_attributes;
-    /// If the return attribute is `kLocation` this stores the location value.
-    std::optional<uint32_t> return_location;
+    /// @returns the function pipeline stage
+    PipelineStage Stage() const { return pipeline_stage_; }
 
-    /// The start target is the first block in a function.
-    Block* start_target = nullptr;
-    /// The end target is the end of the function. It is used as the branch target if a return is
-    /// encountered in the function.
-    FunctionTerminator* end_target = nullptr;
+    /// Sets the workgroup size
+    /// @param x the x size
+    /// @param y the y size
+    /// @param z the z size
+    void SetWorkgroupSize(uint32_t x, uint32_t y, uint32_t z) { workgroup_size_ = {x, y, z}; }
+
+    /// @returns the workgroup size information
+    std::optional<std::array<uint32_t, 3>> WorkgroupSize() const { return workgroup_size_; }
+
+    /// @returns the return type for the function
+    const type::Type* ReturnType() const { return return_type_; }
+
+    /// Sets the return attributes
+    /// @param attrs the attributes to set
+    void SetReturnAttributes(utils::VectorRef<ReturnAttribute> attrs) {
+        return_attributes_ = std::move(attrs);
+    }
+    /// @returns the return attributes
+    utils::VectorRef<ReturnAttribute> ReturnAttributes() const { return return_attributes_; }
+
+    /// Sets the return location
+    /// @param loc the location to set
+    void SetReturnLocation(std::optional<uint32_t> loc) { return_location_ = loc; }
+    /// @returns the return location
+    std::optional<uint32_t> ReturnLocation() const { return return_location_; }
+
+    /// Sets the function parameters
+    /// @param params the function paramters
+    void SetParams(utils::VectorRef<FunctionParam*> params) { params_ = std::move(params); }
+
+    /// @returns the function parameters
+    utils::VectorRef<FunctionParam*> Params() const { return params_; }
+
+    /// Sets the start target for the function
+    /// @param target the start target
+    void SetStartTarget(Block* target) { start_target_ = target; }
+    /// @returns the function start target
+    Block* StartTarget() const { return start_target_; }
+
+    /// Sets the end target for the function
+    /// @param target the end target
+    void SetEndTarget(FunctionTerminator* target) { end_target_ = target; }
+    /// @returns the function end target
+    FunctionTerminator* EndTarget() const { return end_target_; }
+
+  private:
+    Symbol name_;
+    const type::Type* return_type_;
+    PipelineStage pipeline_stage_;
+    std::optional<std::array<uint32_t, 3>> workgroup_size_;
+
+    utils::Vector<ReturnAttribute, 1> return_attributes_;
+    std::optional<uint32_t> return_location_;
+
+    utils::Vector<FunctionParam*, 1> params_;
+
+    Block* start_target_ = nullptr;
+    FunctionTerminator* end_target_ = nullptr;
 };
 
 utils::StringStream& operator<<(utils::StringStream& out, Function::PipelineStage value);
diff --git a/src/tint/ir/function_param.cc b/src/tint/ir/function_param.cc
new file mode 100644
index 0000000..a9570fa
--- /dev/null
+++ b/src/tint/ir/function_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/function_param.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::FunctionParam);
+
+namespace tint::ir {
+
+FunctionParam::FunctionParam(const type::Type* ty) : type_(ty) {}
+
+FunctionParam::~FunctionParam() = default;
+
+}  // namespace tint::ir
diff --git a/src/tint/ir/function_param.h b/src/tint/ir/function_param.h
new file mode 100644
index 0000000..1bbb812
--- /dev/null
+++ b/src/tint/ir/function_param.h
@@ -0,0 +1,46 @@
+// 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_FUNCTION_PARAM_H_
+#define SRC_TINT_IR_FUNCTION_PARAM_H_
+
+#include "src/tint/ir/value.h"
+#include "src/tint/utils/castable.h"
+
+namespace tint::ir {
+
+/// A function parameter in the IR.
+class FunctionParam : public utils::Castable<FunctionParam, Value> {
+  public:
+    /// Constructor
+    /// @param type the type of the var
+    explicit FunctionParam(const type::Type* type);
+    FunctionParam(const FunctionParam& inst) = delete;
+    FunctionParam(FunctionParam&& inst) = delete;
+    ~FunctionParam() override;
+
+    FunctionParam& operator=(const FunctionParam& inst) = delete;
+    FunctionParam& operator=(FunctionParam&& inst) = delete;
+
+    /// @returns the type of the var
+    const type::Type* Type() const override { return type_; }
+
+  private:
+    /// The type of the parameter
+    const type::Type* type_;
+};
+
+}  // namespace tint::ir
+
+#endif  // SRC_TINT_IR_FUNCTION_PARAM_H_
diff --git a/src/tint/ir/if.cc b/src/tint/ir/if.cc
index b59d87f..d235b5e 100644
--- a/src/tint/ir/if.cc
+++ b/src/tint/ir/if.cc
@@ -18,7 +18,7 @@
 
 namespace tint::ir {
 
-If::If(Value* cond) : Base(), condition(cond) {}
+If::If(Value* cond) : Base(), condition_(cond) {}
 
 If::~If() = default;
 
diff --git a/src/tint/ir/if.h b/src/tint/ir/if.h
index aadc5c9..d02c32a 100644
--- a/src/tint/ir/if.h
+++ b/src/tint/ir/if.h
@@ -32,17 +32,36 @@
     /// Constructor
     /// @param cond the if condition
     explicit If(Value* cond);
+    If(const If&) = delete;
+    If(If&&) = delete;
     ~If() override;
 
-    /// The true branch block
+    If& operator=(const If&) = delete;
+    If& operator=(If&&) = delete;
+
+    /// @returns the if condition
+    const Value* Condition() const { return condition_; }
+
+    /// @returns the true branch block
+    const Branch& True() const { return true_; }
+    /// @returns the true branch block
+    Branch& True() { return true_; }
+
+    /// @returns the false branch block
+    const Branch& False() const { return false_; }
+    /// @returns the false branch block
+    Branch& False() { return false_; }
+
+    /// @returns the merge branch block
+    const Branch& Merge() const { return merge_; }
+    /// @returns the merge branch block
+    Branch& Merge() { return merge_; }
+
+  private:
     Branch true_ = {};
-    /// The false branch block
     Branch false_ = {};
-    /// An block to converge the true/false branches. The block always exists, but there maybe no
-    /// branches into it. (e.g. if both branches `return`)
-    Branch merge = {};
-    /// Value holding the condition result
-    const Value* condition = nullptr;
+    Branch merge_ = {};
+    Value* condition_;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/load.cc b/src/tint/ir/load.cc
index 1fe55c0..3b64cdd 100644
--- a/src/tint/ir/load.cc
+++ b/src/tint/ir/load.cc
@@ -19,10 +19,10 @@
 
 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);
+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);
 }
 
 Load::~Load() = default;
diff --git a/src/tint/ir/load.h b/src/tint/ir/load.h
index b15eced..e1a365f 100644
--- a/src/tint/ir/load.h
+++ b/src/tint/ir/load.h
@@ -35,13 +35,14 @@
     Load& operator=(Load&& inst) = delete;
 
     /// @returns the type of the value
-    const type::Type* Type() const override { return result_type; }
+    const type::Type* Type() const override { return result_type_; }
 
-    /// the result type of the instruction
-    const type::Type* result_type = nullptr;
+    /// @returns the avlue being loaded from
+    Value* From() const { return from_; }
 
-    /// the value being loaded
-    Value* from = nullptr;
+  private:
+    const type::Type* result_type_;
+    Value* from_;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/load_test.cc b/src/tint/ir/load_test.cc
index 2c6e5c1..9881e98 100644
--- a/src/tint/ir/load_test.cc
+++ b/src/tint/ir/load_test.cc
@@ -33,12 +33,12 @@
     const auto* inst = b.Load(var);
 
     ASSERT_TRUE(inst->Is<Load>());
-    ASSERT_EQ(inst->from, var);
+    ASSERT_EQ(inst->From(), var);
 
     EXPECT_EQ(inst->Type(), store_type);
 
-    ASSERT_TRUE(inst->from->Is<ir::Var>());
-    EXPECT_EQ(inst->from, var);
+    ASSERT_TRUE(inst->From()->Is<ir::Var>());
+    EXPECT_EQ(inst->From(), var);
 }
 
 TEST_F(IR_InstructionTest, Load_Usage) {
@@ -50,9 +50,9 @@
         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);
+    ASSERT_NE(inst->From(), nullptr);
+    ASSERT_EQ(inst->From()->Usage().Length(), 1u);
+    EXPECT_EQ(inst->From()->Usage()[0], inst);
 }
 
 }  // namespace
diff --git a/src/tint/ir/loop.h b/src/tint/ir/loop.h
index e0066f4..590f49c 100644
--- a/src/tint/ir/loop.h
+++ b/src/tint/ir/loop.h
@@ -26,16 +26,32 @@
   public:
     /// Constructor
     Loop();
+    Loop(const Loop&) = delete;
+    Loop(Loop&&) = delete;
     ~Loop() override;
 
-    /// The start block is the first block in a loop.
-    Branch start = {};
-    /// The continue target of the block.
-    Branch continuing = {};
-    /// The loop merge target. If the `loop` does a `return` then this block may not actually
-    /// end up in the control flow. We need it if the loop does a `break` we know where to break
-    /// too.
-    Branch merge = {};
+    Loop& operator=(const Loop&) = delete;
+    Loop& operator=(Loop&&) = delete;
+
+    /// @returns the switch start branch
+    const Branch& Start() const { return start_; }
+    /// @returns the switch start branch
+    Branch& Start() { return start_; }
+
+    /// @returns the switch continuing branch
+    const Branch& Continuing() const { return continuing_; }
+    /// @returns the switch continuing branch
+    Branch& Continuing() { return continuing_; }
+
+    /// @returns the switch merge branch
+    const Branch& Merge() const { return merge_; }
+    /// @returns the switch merge branch
+    Branch& Merge() { return merge_; }
+
+  private:
+    Branch start_ = {};
+    Branch continuing_ = {};
+    Branch merge_ = {};
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/module.h b/src/tint/ir/module.h
index 244bc0a..c6cf74c 100644
--- a/src/tint/ir/module.h
+++ b/src/tint/ir/module.h
@@ -18,6 +18,7 @@
 #include <string>
 
 #include "src/tint/constant/value.h"
+#include "src/tint/ir/constant.h"
 #include "src/tint/ir/function.h"
 #include "src/tint/ir/instruction.h"
 #include "src/tint/ir/value.h"
@@ -67,7 +68,7 @@
     /// The flow node allocator
     utils::BlockAllocator<FlowNode> flow_nodes;
     /// The constant allocator
-    utils::BlockAllocator<constant::Value> constants;
+    utils::BlockAllocator<constant::Value> constants_arena;
     /// The value allocator
     utils::BlockAllocator<Value> values;
 
@@ -82,6 +83,29 @@
 
     /// The symbol table for the module
     SymbolTable symbols{prog_id_};
+
+    /// ConstantHasher provides a hash function for a constant::Value pointer, hashing the value
+    /// instead of the pointer itself.
+    struct ConstantHasher {
+        /// @param c the constant pointer to create a hash for
+        /// @return the hash value
+        inline std::size_t operator()(const constant::Value* c) const { return c->Hash(); }
+    };
+
+    /// ConstantEquals provides an equality function for two constant::Value pointers, comparing
+    /// their values instead of the pointers.
+    struct ConstantEquals {
+        /// @param a the first constant pointer to compare
+        /// @param b the second constant pointer to compare
+        /// @return the hash value
+        inline bool operator()(const constant::Value* a, const constant::Value* b) const {
+            return a->Equal(b);
+        }
+    };
+
+    /// The map of constant::Value to their ir::Constant.
+    utils::Hashmap<const constant::Value*, ir::Constant*, 16, ConstantHasher, ConstantEquals>
+        constants;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/root_terminator.h b/src/tint/ir/root_terminator.h
index 361aa6d..c6338ba 100644
--- a/src/tint/ir/root_terminator.h
+++ b/src/tint/ir/root_terminator.h
@@ -25,7 +25,12 @@
   public:
     /// Constructor
     RootTerminator();
+    RootTerminator(const RootTerminator&) = delete;
+    RootTerminator(RootTerminator&&) = delete;
     ~RootTerminator() override;
+
+    RootTerminator& operator=(const RootTerminator&) = delete;
+    RootTerminator& operator=(RootTerminator&&) = delete;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/store.cc b/src/tint/ir/store.cc
index 87f3620..8b8de46 100644
--- a/src/tint/ir/store.cc
+++ b/src/tint/ir/store.cc
@@ -19,11 +19,11 @@
 
 namespace tint::ir {
 
-Store::Store(Value* t, Value* f) : Base(), to(t), from(f) {
-    TINT_ASSERT(IR, to);
-    TINT_ASSERT(IR, from);
-    to->AddUsage(this);
-    from->AddUsage(this);
+Store::Store(Value* to, Value* from) : Base(), to_(to), from_(from) {
+    TINT_ASSERT(IR, to_);
+    TINT_ASSERT(IR, from_);
+    to_->AddUsage(this);
+    from_->AddUsage(this);
 }
 
 Store::~Store() = default;
diff --git a/src/tint/ir/store.h b/src/tint/ir/store.h
index 9095c41..af5377d 100644
--- a/src/tint/ir/store.h
+++ b/src/tint/ir/store.h
@@ -34,10 +34,15 @@
     Store& operator=(const Store& inst) = delete;
     Store& operator=(Store&& inst) = delete;
 
-    /// the value being stored to
-    Value* to = nullptr;
-    /// the value being stored
-    Value* from = nullptr;
+    /// @returns the value being stored too
+    Value* To() const { return to_; }
+
+    /// @returns the value being stored
+    Value* From() const { return from_; }
+
+  private:
+    Value* to_;
+    Value* from_;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/store_test.cc b/src/tint/ir/store_test.cc
index 902ca95..1906cc6 100644
--- a/src/tint/ir/store_test.cc
+++ b/src/tint/ir/store_test.cc
@@ -33,10 +33,10 @@
     const auto* inst = b.Store(to, b.Constant(4_i));
 
     ASSERT_TRUE(inst->Is<Store>());
-    ASSERT_EQ(inst->to, to);
+    ASSERT_EQ(inst->To(), to);
 
-    ASSERT_TRUE(inst->from->Is<Constant>());
-    auto lhs = inst->from->As<Constant>()->value;
+    ASSERT_TRUE(inst->From()->Is<Constant>());
+    auto lhs = inst->From()->As<Constant>()->Value();
     ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
@@ -48,13 +48,13 @@
     auto* to = b.Discard();
     const auto* inst = b.Store(to, b.Constant(4_i));
 
-    ASSERT_NE(inst->to, nullptr);
-    ASSERT_EQ(inst->to->Usage().Length(), 1u);
-    EXPECT_EQ(inst->to->Usage()[0], inst);
+    ASSERT_NE(inst->To(), nullptr);
+    ASSERT_EQ(inst->To()->Usage().Length(), 1u);
+    EXPECT_EQ(inst->To()->Usage()[0], inst);
 
-    ASSERT_NE(inst->from, nullptr);
-    ASSERT_EQ(inst->from->Usage().Length(), 1u);
-    EXPECT_EQ(inst->from->Usage()[0], inst);
+    ASSERT_NE(inst->From(), nullptr);
+    ASSERT_EQ(inst->From()->Usage().Length(), 1u);
+    EXPECT_EQ(inst->From()->Usage()[0], inst);
 }
 
 }  // namespace
diff --git a/src/tint/ir/switch.cc b/src/tint/ir/switch.cc
index ad6a145..3bccc83 100644
--- a/src/tint/ir/switch.cc
+++ b/src/tint/ir/switch.cc
@@ -18,7 +18,7 @@
 
 namespace tint::ir {
 
-Switch::Switch(Value* cond) : Base(), condition(cond) {}
+Switch::Switch(Value* cond) : Base(), condition_(cond) {}
 
 Switch::~Switch() = default;
 
diff --git a/src/tint/ir/switch.h b/src/tint/ir/switch.h
index 6be7b62..2e977ec 100644
--- a/src/tint/ir/switch.h
+++ b/src/tint/ir/switch.h
@@ -41,21 +41,40 @@
         utils::Vector<CaseSelector, 4> selectors;
         /// The start block for the case block.
         Branch start = {};
+
+        /// @returns the case start target
+        const Branch& Start() const { return start; }
+        /// @returns the case start target
+        Branch& Start() { return start; }
     };
 
     /// Constructor
     /// @param cond the condition
     explicit Switch(Value* cond);
+    Switch(const Switch&) = delete;
+    Switch(Switch&&) = delete;
     ~Switch() override;
 
-    /// The switch merge target
-    Branch merge = {};
+    Switch& operator=(const Switch&) = delete;
+    Switch& operator=(Switch&&) = delete;
 
-    /// The switch case statements
-    utils::Vector<Case, 4> cases;
+    /// @returns the switch merge branch
+    const Branch& Merge() const { return merge_; }
+    /// @returns the switch merge branch
+    Branch& Merge() { return merge_; }
 
-    /// Value holding the condition result
-    const Value* condition = nullptr;
+    /// @returns the switch cases
+    utils::VectorRef<Case> Cases() const { return cases_; }
+    /// @returns the switch cases
+    utils::Vector<Case, 4>& Cases() { return cases_; }
+
+    /// @returns the condition
+    const Value* Condition() const { return condition_; }
+
+  private:
+    Branch merge_ = {};
+    utils::Vector<Case, 4> cases_;
+    Value* condition_;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/test_helper.h b/src/tint/ir/test_helper.h
index d9055ce..b7279f3 100644
--- a/src/tint/ir/test_helper.h
+++ b/src/tint/ir/test_helper.h
@@ -33,7 +33,6 @@
 class TestHelperBase : public BASE, public ProgramBuilder {
   public:
     TestHelperBase() = default;
-
     ~TestHelperBase() override = default;
 
     /// Build the module, cleaning up the program before returning.
diff --git a/src/tint/ir/to_program.cc b/src/tint/ir/to_program.cc
index b1aa67f..9be189c 100644
--- a/src/tint/ir/to_program.cc
+++ b/src/tint/ir/to_program.cc
@@ -91,14 +91,14 @@
     const ast::Function* Fn(const Function* fn) {
         SCOPED_NESTING();
 
-        auto name = Sym(fn->name);
+        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{};
-        auto ret_ty = Type(fn->return_type);
+        auto ret_ty = Type(fn->ReturnType());
         if (!ret_ty) {
             return nullptr;
         }
-        auto* body = FlowNodeGraph(fn->start_target);
+        auto* body = FlowNodeGraph(fn->StartTarget());
         if (!body) {
             return nullptr;
         }
@@ -126,7 +126,7 @@
                 branch->target,
 
                 [&](const ir::Block* block) {
-                    for (auto* inst : block->instructions) {
+                    for (const auto* inst : block->Instructions()) {
                         auto stmt = Stmt(inst);
                         if (TINT_UNLIKELY(!stmt)) {
                             return kError;
@@ -135,7 +135,7 @@
                             stmts.Push(s);
                         }
                     }
-                    branch = &block->branch;
+                    branch = &block->Branch();
                     return kContinue;
                 },
 
@@ -145,8 +145,8 @@
                         return kError;
                     }
                     stmts.Push(stmt);
-                    branch = &if_->merge;
-                    return branch->target->inbound_branches.IsEmpty() ? kStop : kContinue;
+                    branch = &if_->Merge();
+                    return branch->target->InboundBranches().IsEmpty() ? kStop : kContinue;
                 },
 
                 [&](const ir::Switch* switch_) {
@@ -155,8 +155,8 @@
                         return kError;
                     }
                     stmts.Push(stmt);
-                    branch = &switch_->merge;
-                    return branch->target->inbound_branches.IsEmpty() ? kStop : kContinue;
+                    branch = &switch_->Merge();
+                    return branch->target->InboundBranches().IsEmpty() ? kStop : kContinue;
                 },
 
                 [&](const ir::FunctionTerminator*) {
@@ -189,25 +189,25 @@
     const ast::IfStatement* If(const ir::If* i) {
         SCOPED_NESTING();
 
-        auto* cond = Expr(i->condition);
-        auto* t = FlowNodeGraph(i->true_.target, i->merge.target);
+        auto* cond = Expr(i->Condition());
+        auto* t = FlowNodeGraph(i->True().target, i->Merge().target);
         if (TINT_UNLIKELY(!t)) {
             return nullptr;
         }
 
-        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
+        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));
+            if (auto* else_if = As<ir::If>(NextNonEmptyNode(i->False().target));
                 else_if &&
-                NextNonEmptyNode(i->merge.target) == NextNonEmptyNode(else_if->merge.target)) {
+                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);
+                auto* f = FlowNodeGraph(i->False().target, i->Merge().target);
                 if (!f) {
                     return nullptr;
                 }
@@ -221,16 +221,16 @@
     const ast::SwitchStatement* Switch(const ir::Switch* s) {
         SCOPED_NESTING();
 
-        auto* cond = Expr(s->condition);
+        auto* cond = Expr(s->Condition());
         if (!cond) {
             return nullptr;
         }
 
-        auto cases = utils::Transform(
-            s->cases,  //
+        auto cases = utils::Transform<1>(
+            s->Cases(),  //
             [&](const ir::Switch::Case& c) -> const tint::ast::CaseStatement* {
                 SCOPED_NESTING();
-                auto* body = FlowNodeGraph(c.start.target, s->merge.target);
+                auto* body = FlowNodeGraph(c.start.target, s->Merge().target);
                 if (!body) {
                     return nullptr;
                 }
@@ -292,10 +292,10 @@
     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) {
+                if (!block->Instructions().IsEmpty()) {
                     return false;
                 }
-                node = block->branch.target;
+                node = block->Branch().target;
             } else {
                 return false;
             }
@@ -307,13 +307,13 @@
     const ir::FlowNode* NextNonEmptyNode(const ir::FlowNode* node) {
         while (node) {
             if (auto* block = node->As<ir::Block>()) {
-                for (auto* inst : block->instructions) {
+                for (const auto* inst : block->Instructions()) {
                     // Load instructions will be inlined, so ignore them.
                     if (!inst->Is<ir::Load>()) {
                         return node;
                     }
                 }
-                node = block->branch.target;
+                node = block->Branch().target;
             } else {
                 return node;
             }
@@ -351,8 +351,8 @@
         }
         auto ty = Type(ptr->StoreType());
         const ast::Expression* init = nullptr;
-        if (var->initializer) {
-            init = Expr(var->initializer);
+        if (var->Initializer()) {
+            init = Expr(var->Initializer());
             if (!init) {
                 return nullptr;
             }
@@ -368,18 +368,19 @@
     }
 
     const ast::AssignmentStatement* Store(const ir::Store* store) {
-        auto* expr = Expr(store->from);
-        return b.Assign(NameOf(store->to), expr);
+        auto* expr = Expr(store->From());
+        return b.Assign(NameOf(store->To()), expr);
     }
 
     const ast::CallExpression* Call(const ir::Call* call) {
-        auto args = utils::Transform(call->args, [&](const ir::Value* arg) { return Expr(arg); });
+        auto args =
+            utils::Transform<2>(call->Args(), [&](const ir::Value* arg) { return Expr(arg); });
         if (args.Any(utils::IsNull)) {
             return nullptr;
         }
         return tint::Switch(
             call,  //
-            [&](const ir::UserCall* c) { return b.Call(Sym(c->name), std::move(args)); },
+            [&](const ir::UserCall* c) { return b.Call(Sym(c->Name()), std::move(args)); },
             [&](Default) {
                 UNHANDLED_CASE(call);
                 return nullptr;
@@ -401,18 +402,18 @@
     const ast::Expression* ConstExpr(const ir::Constant* c) {
         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>()); },
-            [&](const type::F32*) { return b.Expr(c->value->ValueAs<f32>()); },
-            [&](const type::F16*) { return b.Expr(c->value->ValueAs<f16>()); },
-            [&](const type::Bool*) { return b.Expr(c->value->ValueAs<bool>()); },
+            [&](const type::I32*) { return b.Expr(c->Value()->ValueAs<i32>()); },
+            [&](const type::U32*) { return b.Expr(c->Value()->ValueAs<u32>()); },
+            [&](const type::F32*) { return b.Expr(c->Value()->ValueAs<f32>()); },
+            [&](const type::F16*) { return b.Expr(c->Value()->ValueAs<f16>()); },
+            [&](const type::Bool*) { return b.Expr(c->Value()->ValueAs<bool>()); },
             [&](Default) {
                 UNHANDLED_CASE(c);
                 return nullptr;
             });
     }
 
-    const ast::Expression* LoadExpr(const ir::Load* l) { return Expr(l->from); }
+    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)); }
 
diff --git a/src/tint/ir/transform/add_empty_entry_point.cc b/src/tint/ir/transform/add_empty_entry_point.cc
index 809a6ad..6788d7c 100644
--- a/src/tint/ir/transform/add_empty_entry_point.cc
+++ b/src/tint/ir/transform/add_empty_entry_point.cc
@@ -29,7 +29,7 @@
 
 void AddEmptyEntryPoint::Run(ir::Module* ir, const DataMap&, DataMap&) const {
     for (auto* func : ir->functions) {
-        if (func->pipeline_stage != Function::PipelineStage::kUndefined) {
+        if (func->Stage() != Function::PipelineStage::kUndefined) {
             return;
         }
     }
@@ -38,7 +38,7 @@
     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);
+    ep->StartTarget()->BranchTo(ep->EndTarget());
     ir->functions.Push(ep);
 }
 
diff --git a/src/tint/ir/transform/add_empty_entry_point_test.cc b/src/tint/ir/transform/add_empty_entry_point_test.cc
index baba8e0..a363d83 100644
--- a/src/tint/ir/transform/add_empty_entry_point_test.cc
+++ b/src/tint/ir/transform/add_empty_entry_point_test.cc
@@ -40,7 +40,7 @@
 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);
+    ep->StartTarget()->BranchTo(ep->EndTarget());
     mod.functions.Push(ep);
 
     auto* expect = R"(
diff --git a/src/tint/ir/unary.cc b/src/tint/ir/unary.cc
index db13d5d..2a0ee94 100644
--- a/src/tint/ir/unary.cc
+++ b/src/tint/ir/unary.cc
@@ -19,8 +19,8 @@
 
 namespace tint::ir {
 
-Unary::Unary(Kind k, const type::Type* res_ty, Value* val)
-    : kind(k), result_type(res_ty), val_(val) {
+Unary::Unary(enum Kind k, const type::Type* res_ty, Value* val)
+    : kind_(k), result_type_(res_ty), val_(val) {
     TINT_ASSERT(IR, val_);
     val_->AddUsage(this);
 }
diff --git a/src/tint/ir/unary.h b/src/tint/ir/unary.h
index 46edd45..98ad9b8 100644
--- a/src/tint/ir/unary.h
+++ b/src/tint/ir/unary.h
@@ -33,7 +33,7 @@
     /// @param kind the kind of unary instruction
     /// @param result_type the result type
     /// @param val the input value for the instruction
-    Unary(Kind kind, const type::Type* result_type, Value* val);
+    Unary(enum Kind kind, const type::Type* result_type, Value* val);
     Unary(const Unary& inst) = delete;
     Unary(Unary&& inst) = delete;
     ~Unary() override;
@@ -42,19 +42,18 @@
     Unary& operator=(Unary&& inst) = delete;
 
     /// @returns the type of the value
-    const type::Type* Type() const override { return result_type; }
+    const type::Type* Type() const override { return result_type_; }
 
     /// @returns the value for the instruction
     const Value* Val() const { return val_; }
 
-    /// the kind of unary instruction
-    Kind kind = Kind::kNegation;
-
-    /// the result type of the instruction
-    const type::Type* result_type = nullptr;
+    /// @returns the kind of unary instruction
+    enum Kind Kind() const { return kind_; }
 
   private:
-    Value* val_ = nullptr;
+    enum Kind kind_;
+    const type::Type* result_type_;
+    Value* val_;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/unary_test.cc b/src/tint/ir/unary_test.cc
index 280c7de..6cefed3 100644
--- a/src/tint/ir/unary_test.cc
+++ b/src/tint/ir/unary_test.cc
@@ -29,10 +29,10 @@
     const auto* inst = b.Complement(b.ir.types.Get<type::I32>(), b.Constant(4_i));
 
     ASSERT_TRUE(inst->Is<Unary>());
-    EXPECT_EQ(inst->kind, Unary::Kind::kComplement);
+    EXPECT_EQ(inst->Kind(), Unary::Kind::kComplement);
 
     ASSERT_TRUE(inst->Val()->Is<Constant>());
-    auto lhs = inst->Val()->As<Constant>()->value;
+    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>());
 }
@@ -43,10 +43,10 @@
     const auto* inst = b.Negation(b.ir.types.Get<type::I32>(), b.Constant(4_i));
 
     ASSERT_TRUE(inst->Is<Unary>());
-    EXPECT_EQ(inst->kind, Unary::Kind::kNegation);
+    EXPECT_EQ(inst->Kind(), Unary::Kind::kNegation);
 
     ASSERT_TRUE(inst->Val()->Is<Constant>());
-    auto lhs = inst->Val()->As<Constant>()->value;
+    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>());
 }
@@ -56,7 +56,7 @@
     Builder b{mod};
     const auto* inst = b.Negation(b.ir.types.Get<type::I32>(), b.Constant(4_i));
 
-    EXPECT_EQ(inst->kind, Unary::Kind::kNegation);
+    EXPECT_EQ(inst->Kind(), Unary::Kind::kNegation);
 
     ASSERT_NE(inst->Val(), nullptr);
     ASSERT_EQ(inst->Val()->Usage().Length(), 1u);
diff --git a/src/tint/ir/user_call.cc b/src/tint/ir/user_call.cc
index f718284..e44e285 100644
--- a/src/tint/ir/user_call.cc
+++ b/src/tint/ir/user_call.cc
@@ -23,7 +23,7 @@
 namespace tint::ir {
 
 UserCall::UserCall(const type::Type* ty, Symbol n, utils::VectorRef<Value*> arguments)
-    : Base(ty, std::move(arguments)), name(n) {}
+    : Base(ty, std::move(arguments)), name_(n) {}
 
 UserCall::~UserCall() = default;
 
diff --git a/src/tint/ir/user_call.h b/src/tint/ir/user_call.h
index 5ada8f8..ba52e20 100644
--- a/src/tint/ir/user_call.h
+++ b/src/tint/ir/user_call.h
@@ -36,8 +36,11 @@
     UserCall& operator=(const UserCall& inst) = delete;
     UserCall& operator=(UserCall&& inst) = delete;
 
-    /// The function name
-    Symbol name;
+    /// @returns the called function name
+    Symbol Name() const { return name_; }
+
+  private:
+    Symbol name_;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/var.cc b/src/tint/ir/var.cc
index 9bf4b8d..9b43329 100644
--- a/src/tint/ir/var.cc
+++ b/src/tint/ir/var.cc
@@ -19,8 +19,14 @@
 
 namespace tint::ir {
 
-Var::Var(const type::Type* ty) : type(ty) {}
+Var::Var(const type::Type* ty) : type_(ty) {}
 
 Var::~Var() = default;
 
+void Var::SetInitializer(Value* initializer) {
+    initializer_ = initializer;
+    initializer_->AddUsage(this);
+    // TODO(dsinclair): Probably should do a RemoveUsage on an existing initializer if set
+}
+
 }  // namespace tint::ir
diff --git a/src/tint/ir/var.h b/src/tint/ir/var.h
index c874a62..67ccf0d 100644
--- a/src/tint/ir/var.h
+++ b/src/tint/ir/var.h
@@ -36,13 +36,17 @@
     Var& operator=(Var&& inst) = delete;
 
     /// @returns the type of the var
-    const type::Type* Type() const override { return type; }
+    const type::Type* Type() const override { return type_; }
 
-    /// the result type of the instruction
-    const type::Type* type = nullptr;
+    /// Sets the var initializer
+    /// @param initializer the initializer
+    void SetInitializer(Value* initializer);
+    /// @returns the initializer
+    const Value* Initializer() const { return initializer_; }
 
-    /// The optional initializer
-    Value* initializer = nullptr;
+  private:
+    const type::Type* type_;
+    Value* initializer_ = nullptr;
 };
 
 }  // namespace tint::ir
diff --git a/src/tint/transform/manager_test.cc b/src/tint/transform/manager_test.cc
index 3b0ee25..a81f7bf 100644
--- a/src/tint/transform/manager_test.cc
+++ b/src/tint/transform/manager_test.cc
@@ -52,7 +52,7 @@
         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);
+        func->StartTarget()->BranchTo(func->EndTarget());
         mod->functions.Push(func);
     }
 };
@@ -70,7 +70,7 @@
     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);
+    func->StartTarget()->BranchTo(func->EndTarget());
     builder.ir.functions.Push(func);
     return mod;
 }
@@ -102,9 +102,10 @@
     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");
+    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) {
@@ -149,9 +150,9 @@
 
     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");
+    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) {
@@ -164,9 +165,9 @@
 
     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");
+    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
 
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.cc b/src/tint/writer/spirv/ir/generator_impl_ir.cc
index 4ec4d9d..bd9a735 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.cc
@@ -19,6 +19,7 @@
 #include "src/tint/ir/block.h"
 #include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
+#include "src/tint/ir/load.h"
 #include "src/tint/ir/module.h"
 #include "src/tint/ir/store.h"
 #include "src/tint/ir/transform/add_empty_entry_point.h"
@@ -109,7 +110,7 @@
 }
 
 uint32_t GeneratorImplIr::Constant(const ir::Constant* constant) {
-    return Constant(constant->value);
+    return Constant(constant->Value());
 }
 
 uint32_t GeneratorImplIr::Constant(const constant::Value* constant) {
@@ -213,15 +214,15 @@
     auto id = module_.NextId();
 
     // Emit the function name.
-    module_.PushDebug(spv::Op::OpName, {id, Operand(func->name.Name())});
+    module_.PushDebug(spv::Op::OpName, {id, Operand(func->Name().Name())});
 
     // Emit OpEntryPoint and OpExecutionMode declarations if needed.
-    if (func->pipeline_stage != ir::Function::PipelineStage::kUndefined) {
+    if (func->Stage() != ir::Function::PipelineStage::kUndefined) {
         EmitEntryPoint(func, id);
     }
 
     // Get the ID for the return type.
-    auto return_type_id = Type(func->return_type);
+    auto return_type_id = Type(func->ReturnType());
 
     // Get the ID for the function type (creating it if needed).
     // TODO(jrprice): Add the parameter types when they are supported in the IR.
@@ -247,7 +248,7 @@
     TINT_DEFER(current_function_ = Function());
 
     // Emit the body of the function.
-    EmitBlock(func->start_target);
+    EmitBlock(func->StartTarget());
 
     // Add the function to the module.
     module_.PushFunction(current_function_);
@@ -255,13 +256,13 @@
 
 void GeneratorImplIr::EmitEntryPoint(const ir::Function* func, uint32_t id) {
     SpvExecutionModel stage = SpvExecutionModelMax;
-    switch (func->pipeline_stage) {
+    switch (func->Stage()) {
         case ir::Function::PipelineStage::kCompute: {
             stage = SpvExecutionModelGLCompute;
             module_.PushExecutionMode(
                 spv::Op::OpExecutionMode,
-                {id, U32Operand(SpvExecutionModeLocalSize), func->workgroup_size->at(0),
-                 func->workgroup_size->at(1), func->workgroup_size->at(2)});
+                {id, U32Operand(SpvExecutionModeLocalSize), func->WorkgroupSize()->at(0),
+                 func->WorkgroupSize()->at(1), func->WorkgroupSize()->at(2)});
             break;
         }
         case ir::Function::PipelineStage::kFragment: {
@@ -281,7 +282,7 @@
     }
 
     // TODO(jrprice): Add the interface list of all referenced global variables.
-    module_.PushEntryPoint(spv::Op::OpEntryPoint, {U32Operand(stage), id, func->name.Name()});
+    module_.PushEntryPoint(spv::Op::OpEntryPoint, {U32Operand(stage), id, func->Name().Name()});
 }
 
 void GeneratorImplIr::EmitBlock(const ir::Block* block) {
@@ -292,10 +293,11 @@
     }
 
     // Emit the instructions.
-    for (auto* inst : block->instructions) {
+    for (const auto* inst : block->Instructions()) {
         auto result = Switch(
             inst,  //
             [&](const ir::Binary* b) { return EmitBinary(b); },
+            [&](const ir::Load* l) { return EmitLoad(l); },
             [&](const ir::Store* s) {
                 EmitStore(s);
                 return 0u;
@@ -311,43 +313,43 @@
 
     // Handle the branch at the end of the block.
     Switch(
-        block->branch.target,
+        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()) {
+            if (!block->Branch().args.IsEmpty()) {
                 TINT_ICE(Writer, diagnostics_) << "unimplemented return value";
             }
             current_function_.push_inst(spv::Op::OpReturn, {});
         },
         [&](Default) {
-            if (!block->branch.target) {
+            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;
+                    << "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>();
+    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.
+    // 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) {
+    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) {
+    if (!false_block->Instructions().IsEmpty() || false_block->Branch().target != merge_block) {
         false_label = Label(false_block);
     }
 
@@ -355,7 +357,7 @@
     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});
+                                {Value(i->Condition()), true_label, false_label});
 
     // Emit the `true` and `false` blocks, if they're not being skipped.
     if (true_label != merge_label) {
@@ -374,7 +376,7 @@
 
     // Determine the opcode.
     spv::Op op = spv::Op::Max;
-    switch (binary->kind) {
+    switch (binary->Kind()) {
         case ir::Binary::Kind::kAdd: {
             op = binary->Type()->is_integer_scalar_or_vector() ? spv::Op::OpIAdd : spv::Op::OpFAdd;
             break;
@@ -385,7 +387,7 @@
         }
         default: {
             TINT_ICE(Writer, diagnostics_)
-                << "unimplemented binary instruction: " << static_cast<uint32_t>(binary->kind);
+                << "unimplemented binary instruction: " << static_cast<uint32_t>(binary->Kind());
         }
     }
 
@@ -396,8 +398,14 @@
     return id;
 }
 
+uint32_t GeneratorImplIr::EmitLoad(const ir::Load* load) {
+    auto id = module_.NextId();
+    current_function_.push_inst(spv::Op::OpLoad, {Type(load->Type()), id, Value(load->From())});
+    return id;
+}
+
 void GeneratorImplIr::EmitStore(const ir::Store* store) {
-    current_function_.push_inst(spv::Op::OpStore, {Value(store->to), Value(store->from)});
+    current_function_.push_inst(spv::Op::OpStore, {Value(store->To()), Value(store->From())});
 }
 
 uint32_t GeneratorImplIr::EmitVar(const ir::Var* var) {
@@ -409,8 +417,8 @@
     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)});
+        if (var->Initializer()) {
+            current_function_.push_inst(spv::Op::OpStore, {id, Value(var->Initializer())});
         }
     } else {
         TINT_ICE(Writer, diagnostics_)
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.h b/src/tint/writer/spirv/ir/generator_impl_ir.h
index 66ffe48..b72201f 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir.h
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.h
@@ -32,6 +32,7 @@
 class Block;
 class If;
 class Function;
+class Load;
 class Module;
 class Store;
 class Value;
@@ -106,6 +107,11 @@
     /// @returns the result ID of the instruction
     uint32_t EmitBinary(const ir::Binary* binary);
 
+    /// Emit a load instruction.
+    /// @param load the load instruction to emit
+    /// @returns the result ID of the instruction
+    uint32_t EmitLoad(const ir::Load* load);
+
     /// Emit a store instruction.
     /// @param store the store instruction to emit
     void EmitStore(const ir::Store* store);
@@ -151,25 +157,6 @@
         }
     };
 
-    /// ConstantHasher provides a hash function for a constant::Value pointer, hashing the value
-    /// instead of the pointer itself.
-    struct ConstantHasher {
-        /// @param c the constant::Value pointer to create a hash for
-        /// @return the hash value
-        inline std::size_t operator()(const constant::Value* c) const { return c->Hash(); }
-    };
-
-    /// ConstantEquals provides an equality function for two constant::Value pointers, comparing
-    /// their values instead of the pointers.
-    struct ConstantEquals {
-        /// @param a the first constant::Value pointer to compare
-        /// @param b the second constant::Value pointer to compare
-        /// @return the hash value
-        inline bool operator()(const constant::Value* a, const constant::Value* b) const {
-            return a->Equal(b);
-        }
-    };
-
     /// The map of types to their result IDs.
     utils::Hashmap<const type::Type*, uint32_t, 8> types_;
 
@@ -177,7 +164,7 @@
     utils::Hashmap<FunctionType, uint32_t, 8, FunctionType::Hasher> function_types_;
 
     /// The map of constants to their result IDs.
-    utils::Hashmap<const constant::Value*, uint32_t, 16, ConstantHasher, ConstantEquals> constants_;
+    utils::Hashmap<const constant::Value*, uint32_t, 16> constants_;
 
     /// The map of instructions to their result IDs.
     utils::Hashmap<const ir::Instruction*, uint32_t, 8> instructions_;
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
index 30dc059..e9231f3 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
@@ -21,10 +21,10 @@
 
 TEST_F(SpvGeneratorImplTest, Binary_Add_I32) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
-    func->start_target->instructions.Push(
-        b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i)));
+    func->StartTarget()->SetInstructions(
+        utils::Vector{b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i))});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -43,10 +43,10 @@
 
 TEST_F(SpvGeneratorImplTest, Binary_Add_U32) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
-    func->start_target->instructions.Push(
-        b.Add(mod.types.Get<type::U32>(), b.Constant(1_u), b.Constant(2_u)));
+    func->StartTarget()->SetInstructions(
+        utils::Vector{b.Add(mod.types.Get<type::U32>(), b.Constant(1_u), b.Constant(2_u))});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -65,10 +65,10 @@
 
 TEST_F(SpvGeneratorImplTest, Binary_Add_F32) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
-    func->start_target->instructions.Push(
-        b.Add(mod.types.Get<type::F32>(), b.Constant(1_f), b.Constant(2_f)));
+    func->StartTarget()->SetInstructions(
+        utils::Vector{b.Add(mod.types.Get<type::F32>(), b.Constant(1_f), b.Constant(2_f))});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -87,10 +87,10 @@
 
 TEST_F(SpvGeneratorImplTest, Binary_Sub_I32) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
-    func->start_target->instructions.Push(
-        b.Subtract(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i)));
+    func->StartTarget()->SetInstructions(
+        utils::Vector{b.Subtract(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i))});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -109,10 +109,10 @@
 
 TEST_F(SpvGeneratorImplTest, Binary_Sub_U32) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
-    func->start_target->instructions.Push(
-        b.Subtract(mod.types.Get<type::U32>(), b.Constant(1_u), b.Constant(2_u)));
+    func->StartTarget()->SetInstructions(
+        utils::Vector{b.Subtract(mod.types.Get<type::U32>(), b.Constant(1_u), b.Constant(2_u))});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -131,10 +131,10 @@
 
 TEST_F(SpvGeneratorImplTest, Binary_Sub_F32) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
-    func->start_target->instructions.Push(
-        b.Subtract(mod.types.Get<type::F32>(), b.Constant(1_f), b.Constant(2_f)));
+    func->StartTarget()->SetInstructions(
+        utils::Vector{b.Subtract(mod.types.Get<type::F32>(), b.Constant(1_f), b.Constant(2_f))});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -153,17 +153,17 @@
 
 TEST_F(SpvGeneratorImplTest, Binary_Sub_Vec2i) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
-    auto* lhs = mod.constants.Create<constant::Composite>(
+    auto* lhs = mod.constants_arena.Create<constant::Composite>(
         mod.types.Get<type::Vector>(mod.types.Get<type::I32>(), 2u),
-        utils::Vector{b.Constant(42_i)->value, b.Constant(-1_i)->value}, false, false);
-    auto* rhs = mod.constants.Create<constant::Composite>(
+        utils::Vector{b.Constant(42_i)->Value(), b.Constant(-1_i)->Value()}, false, false);
+    auto* rhs = mod.constants_arena.Create<constant::Composite>(
         mod.types.Get<type::Vector>(mod.types.Get<type::I32>(), 2u),
-        utils::Vector{b.Constant(0_i)->value, b.Constant(-43_i)->value}, false, false);
-    func->start_target->instructions.Push(
-        b.Subtract(mod.types.Get<type::Vector>(mod.types.Get<type::I32>(), 2u), b.Constant(lhs),
-                   b.Constant(rhs)));
+        utils::Vector{b.Constant(0_i)->Value(), b.Constant(-43_i)->Value()}, false, false);
+    func->StartTarget()->SetInstructions(
+        utils::Vector{b.Subtract(mod.types.Get<type::Vector>(mod.types.Get<type::I32>(), 2u),
+                                 b.Constant(lhs), b.Constant(rhs))});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -187,21 +187,21 @@
 
 TEST_F(SpvGeneratorImplTest, Binary_Sub_Vec4f) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
-    auto* lhs = mod.constants.Create<constant::Composite>(
+    auto* lhs = mod.constants_arena.Create<constant::Composite>(
         mod.types.Get<type::Vector>(mod.types.Get<type::F32>(), 4u),
-        utils::Vector{b.Constant(42_f)->value, b.Constant(-1_f)->value, b.Constant(0_f)->value,
-                      b.Constant(1.25_f)->value},
+        utils::Vector{b.Constant(42_f)->Value(), b.Constant(-1_f)->Value(),
+                      b.Constant(0_f)->Value(), b.Constant(1.25_f)->Value()},
         false, false);
-    auto* rhs = mod.constants.Create<constant::Composite>(
+    auto* rhs = mod.constants_arena.Create<constant::Composite>(
         mod.types.Get<type::Vector>(mod.types.Get<type::F32>(), 4u),
-        utils::Vector{b.Constant(0_f)->value, b.Constant(1.25_f)->value, b.Constant(-42_f)->value,
-                      b.Constant(1_f)->value},
+        utils::Vector{b.Constant(0_f)->Value(), b.Constant(1.25_f)->Value(),
+                      b.Constant(-42_f)->Value(), b.Constant(1_f)->Value()},
         false, false);
-    func->start_target->instructions.Push(
-        b.Subtract(mod.types.Get<type::Vector>(mod.types.Get<type::F32>(), 4u), b.Constant(lhs),
-                   b.Constant(rhs)));
+    func->StartTarget()->SetInstructions(
+        utils::Vector{b.Subtract(mod.types.Get<type::Vector>(mod.types.Get<type::F32>(), 4u),
+                                 b.Constant(lhs), b.Constant(rhs))});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -227,11 +227,10 @@
 
 TEST_F(SpvGeneratorImplTest, Binary_Chain) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
     auto* a = b.Subtract(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i));
-    func->start_target->instructions.Push(a);
-    func->start_target->instructions.Push(b.Add(mod.types.Get<type::I32>(), a, a));
+    func->StartTarget()->SetInstructions(utils::Vector{a, b.Add(mod.types.Get<type::I32>(), a, a)});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc
index 95fce03..6ab48aa 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc
@@ -65,9 +65,9 @@
 TEST_F(SpvGeneratorImplTest, Constant_Vec4Bool) {
     auto* t = b.Constant(true);
     auto* f = b.Constant(false);
-    auto* v = mod.constants.Create<constant::Composite>(
+    auto* v = mod.constants_arena.Create<constant::Composite>(
         mod.types.Get<type::Vector>(mod.types.Get<type::Bool>(), 4u),
-        utils::Vector{t->value, f->value, f->value, t->value}, false, true);
+        utils::Vector{t->Value(), f->Value(), f->Value(), t->Value()}, false, true);
     generator_.Constant(b.Constant(v));
     EXPECT_EQ(DumpTypes(), R"(%3 = OpTypeBool
 %2 = OpTypeVector %3 4
@@ -81,8 +81,9 @@
     auto* i = mod.types.Get<type::I32>();
     auto* i_42 = b.Constant(i32(42));
     auto* i_n1 = b.Constant(i32(-1));
-    auto* v = mod.constants.Create<constant::Composite>(
-        mod.types.Get<type::Vector>(i, 2u), utils::Vector{i_42->value, i_n1->value}, false, false);
+    auto* v = mod.constants_arena.Create<constant::Composite>(
+        mod.types.Get<type::Vector>(i, 2u), utils::Vector{i_42->Value(), i_n1->Value()}, false,
+        false);
     generator_.Constant(b.Constant(v));
     EXPECT_EQ(DumpTypes(), R"(%3 = OpTypeInt 32 1
 %2 = OpTypeVector %3 2
@@ -97,9 +98,9 @@
     auto* u_42 = b.Constant(u32(42));
     auto* u_0 = b.Constant(u32(0));
     auto* u_4b = b.Constant(u32(4000000000));
-    auto* v = mod.constants.Create<constant::Composite>(
-        mod.types.Get<type::Vector>(u, 3u), utils::Vector{u_42->value, u_0->value, u_4b->value},
-        false, true);
+    auto* v = mod.constants_arena.Create<constant::Composite>(
+        mod.types.Get<type::Vector>(u, 3u),
+        utils::Vector{u_42->Value(), u_0->Value(), u_4b->Value()}, false, true);
     generator_.Constant(b.Constant(v));
     EXPECT_EQ(DumpTypes(), R"(%3 = OpTypeInt 32 0
 %2 = OpTypeVector %3 3
@@ -116,9 +117,9 @@
     auto* f_0 = b.Constant(f32(0));
     auto* f_q = b.Constant(f32(0.25));
     auto* f_n1 = b.Constant(f32(-1));
-    auto* v = mod.constants.Create<constant::Composite>(
+    auto* v = mod.constants_arena.Create<constant::Composite>(
         mod.types.Get<type::Vector>(f, 4u),
-        utils::Vector{f_42->value, f_0->value, f_q->value, f_n1->value}, false, true);
+        utils::Vector{f_42->Value(), f_0->Value(), f_q->Value(), f_n1->Value()}, false, true);
     generator_.Constant(b.Constant(v));
     EXPECT_EQ(DumpTypes(), R"(%3 = OpTypeFloat 32
 %2 = OpTypeVector %3 4
@@ -134,8 +135,9 @@
     auto* h = mod.types.Get<type::F16>();
     auto* h_42 = b.Constant(f16(42));
     auto* h_q = b.Constant(f16(0.25));
-    auto* v = mod.constants.Create<constant::Composite>(
-        mod.types.Get<type::Vector>(h, 2u), utils::Vector{h_42->value, h_q->value}, false, false);
+    auto* v = mod.constants_arena.Create<constant::Composite>(
+        mod.types.Get<type::Vector>(h, 2u), utils::Vector{h_42->Value(), h_q->Value()}, false,
+        false);
     generator_.Constant(b.Constant(v));
     EXPECT_EQ(DumpTypes(), R"(%3 = OpTypeFloat 16
 %2 = OpTypeVector %3 2
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc
index 77b4a62..d2af246 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc
@@ -19,7 +19,7 @@
 
 TEST_F(SpvGeneratorImplTest, Function_Empty) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -35,7 +35,7 @@
 // Test that we do not emit the same function type more than once.
 TEST_F(SpvGeneratorImplTest, Function_DeduplicateType) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
     generator_.EmitFunction(func);
     generator_.EmitFunction(func);
@@ -48,7 +48,7 @@
 TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Compute) {
     auto* func = b.CreateFunction(mod.symbols.Register("main"), mod.types.Get<type::Void>(),
                                   ir::Function::PipelineStage::kCompute, {{32, 4, 1}});
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint GLCompute %1 "main"
@@ -66,7 +66,7 @@
 TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Fragment) {
     auto* func = b.CreateFunction(mod.symbols.Register("main"), mod.types.Get<type::Void>(),
                                   ir::Function::PipelineStage::kFragment);
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Fragment %1 "main"
@@ -84,7 +84,7 @@
 TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Vertex) {
     auto* func = b.CreateFunction(mod.symbols.Register("main"), mod.types.Get<type::Void>(),
                                   ir::Function::PipelineStage::kVertex);
-    b.Branch(func->start_target, func->end_target);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Vertex %1 "main"
@@ -101,15 +101,15 @@
 TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Multiple) {
     auto* f1 = b.CreateFunction(mod.symbols.Register("main1"), mod.types.Get<type::Void>(),
                                 ir::Function::PipelineStage::kCompute, {{32, 4, 1}});
-    b.Branch(f1->start_target, f1->end_target);
+    f1->StartTarget()->BranchTo(f1->EndTarget());
 
     auto* f2 = b.CreateFunction(mod.symbols.Register("main2"), mod.types.Get<type::Void>(),
                                 ir::Function::PipelineStage::kCompute, {{8, 2, 16}});
-    b.Branch(f2->start_target, f2->end_target);
+    f2->StartTarget()->BranchTo(f2->EndTarget());
 
     auto* f3 = b.CreateFunction(mod.symbols.Register("main3"), mod.types.Get<type::Void>(),
                                 ir::Function::PipelineStage::kFragment);
-    b.Branch(f3->start_target, f3->end_target);
+    f3->StartTarget()->BranchTo(f3->EndTarget());
 
     generator_.EmitFunction(f1);
     generator_.EmitFunction(f2);
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
index 7b41184..fa8cc6b 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
@@ -23,11 +23,11 @@
     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);
+    i->True().target->As<ir::Block>()->BranchTo(i->Merge().target);
+    i->False().target->As<ir::Block>()->BranchTo(i->Merge().target);
+    i->Merge().target->As<ir::Block>()->BranchTo(func->EndTarget());
 
-    b.Branch(func->start_target, i);
+    func->StartTarget()->BranchTo(i);
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -49,15 +49,15 @@
     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);
+    i->False().target->As<ir::Block>()->BranchTo(i->Merge().target);
+    i->Merge().target->As<ir::Block>()->BranchTo(func->EndTarget());
 
-    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);
+    auto* true_block = i->True().target->As<ir::Block>();
+    true_block->SetInstructions(
+        utils::Vector{b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(1_i))});
+    true_block->BranchTo(i->Merge().target);
 
-    b.Branch(func->start_target, i);
+    func->StartTarget()->BranchTo(i);
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -84,15 +84,15 @@
     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);
+    i->True().target->As<ir::Block>()->BranchTo(i->Merge().target);
+    i->Merge().target->As<ir::Block>()->BranchTo(func->EndTarget());
 
-    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);
+    auto* false_block = i->False().target->As<ir::Block>();
+    false_block->SetInstructions(
+        utils::Vector{b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(1_i))});
+    false_block->BranchTo(i->Merge().target);
 
-    b.Branch(func->start_target, i);
+    func->StartTarget()->BranchTo(i);
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -119,11 +119,11 @@
     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;
+    i->True().target->As<ir::Block>()->BranchTo(func->EndTarget());
+    i->False().target->As<ir::Block>()->BranchTo(func->EndTarget());
+    i->Merge().target->As<ir::Block>()->BranchTo(nullptr);
 
-    b.Branch(func->start_target, i);
+    func->StartTarget()->BranchTo(i);
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
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
index ff94862..20e5167 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
@@ -22,12 +22,11 @@
 
 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);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
     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->StartTarget()->SetInstructions(utils::Vector{b.Declare(ty)});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -45,13 +44,14 @@
 
 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);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
     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);
+    v->SetInitializer(b.Constant(42_i));
+
+    func->StartTarget()->SetInstructions(utils::Vector{v});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -71,12 +71,12 @@
 
 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);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
     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->StartTarget()->SetInstructions(utils::Vector{v});
     mod.SetName(v, "myvar");
 
     generator_.EmitFunction(func);
@@ -96,22 +96,22 @@
 
 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);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
     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);
+    v->SetInitializer(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);
+    i->False().target->As<ir::Block>()->BranchTo(func->EndTarget());
+    i->Merge().target->As<ir::Block>()->BranchTo(func->EndTarget());
 
-    auto* true_block = i->true_.target->As<ir::Block>();
-    true_block->instructions.Push(v);
-    b.Branch(true_block, i->merge.target);
+    auto* true_block = i->True().target->As<ir::Block>();
+    true_block->SetInstructions(utils::Vector{v});
+    true_block->BranchTo(i->Merge().target);
 
-    b.Branch(func->start_target, i);
+    func->StartTarget()->BranchTo(i);
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -138,15 +138,39 @@
 )");
 }
 
+TEST_F(SpvGeneratorImplTest, FunctionVar_Load) {
+    auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
+    func->StartTarget()->BranchTo(func->EndTarget());
+
+    auto* store_ty = mod.types.Get<type::I32>();
+    auto* ty = mod.types.Get<type::Pointer>(store_ty, builtin::AddressSpace::kFunction,
+                                            builtin::Access::kReadWrite);
+    auto* v = b.Declare(ty);
+    func->StartTarget()->SetInstructions(utils::Vector{v, b.Load(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
+%8 = OpLoad %7 %5
+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);
+    func->StartTarget()->BranchTo(func->EndTarget());
 
     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)));
+    func->StartTarget()->SetInstructions(utils::Vector{v, b.Store(v, b.Constant(42_i))});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"