diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index ea16d39eea..36311088b1 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -185,18 +185,10 @@ Unary* Builder::CreateUnary(Unary::Kind kind, const type::Type* type, Value* val
     return ir.values.Create<ir::Unary>(kind, type, val);
 }
 
-Unary* Builder::AddressOf(const type::Type* type, Value* val) {
-    return CreateUnary(Unary::Kind::kAddressOf, type, val);
-}
-
 Unary* Builder::Complement(const type::Type* type, Value* val) {
     return CreateUnary(Unary::Kind::kComplement, type, val);
 }
 
-Unary* Builder::Indirection(const type::Type* type, Value* val) {
-    return CreateUnary(Unary::Kind::kIndirection, type, val);
-}
-
 Unary* Builder::Negation(const type::Type* type, Value* val) {
     return CreateUnary(Unary::Kind::kNegation, type, val);
 }
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index e94440435e..f4bc2c5008 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -279,24 +279,12 @@ class Builder {
     /// @returns the operation
     Unary* CreateUnary(Unary::Kind kind, const type::Type* type, Value* val);
 
-    /// Creates an AddressOf operation
-    /// @param type the result type of the expression
-    /// @param val the value
-    /// @returns the operation
-    Unary* AddressOf(const type::Type* type, Value* val);
-
     /// Creates a Complement operation
     /// @param type the result type of the expression
     /// @param val the value
     /// @returns the operation
     Unary* Complement(const type::Type* type, Value* val);
 
-    /// Creates an Indirection operation
-    /// @param type the result type of the expression
-    /// @param val the value
-    /// @returns the operation
-    Unary* Indirection(const type::Type* type, Value* val);
-
     /// Creates a Negation operation
     /// @param type the result type of the expression
     /// @param val the value
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 52e038b01f..ed5cc259e5 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -530,15 +530,9 @@ void Disassembler::EmitUnary(const Unary* u) {
     EmitValue(u);
     out_ << " = ";
     switch (u->kind) {
-        case Unary::Kind::kAddressOf:
-            out_ << "addr_of";
-            break;
         case Unary::Kind::kComplement:
             out_ << "complement";
             break;
-        case Unary::Kind::kIndirection:
-            out_ << "indirection";
-            break;
         case Unary::Kind::kNegation:
             out_ << "negation";
             break;
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index d42ea0bdb9..62fd3b9833 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -82,6 +82,7 @@
 #include "src/tint/sem/value_expression.h"
 #include "src/tint/sem/variable.h"
 #include "src/tint/switch.h"
+#include "src/tint/type/pointer.h"
 #include "src/tint/type/reference.h"
 #include "src/tint/type/void.h"
 #include "src/tint/utils/defer.h"
@@ -853,7 +854,11 @@ class Impl {
         return tint::Switch(  //
             var,
             [&](const ast::Var* v) {
-                auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx);
+                auto* ref = sem->Type()->As<type::Reference>();
+                auto* ty = builder_.ir.types.Get<type::Pointer>(
+                    ref->StoreType()->Clone(clone_ctx_.type_ctx), ref->AddressSpace(),
+                    ref->Access());
+
                 auto* val = builder_.Declare(ty);
                 current_flow_block_->instructions.Push(val);
 
@@ -915,14 +920,12 @@ class Impl {
         Instruction* inst = nullptr;
         switch (expr->op) {
             case ast::UnaryOp::kAddressOf:
-                inst = builder_.AddressOf(ty, val.Get());
-                break;
+            case ast::UnaryOp::kIndirection:
+                // 'address-of' and 'indirection' just fold away and we propagate the pointer.
+                return val;
             case ast::UnaryOp::kComplement:
                 inst = builder_.Complement(ty, val.Get());
                 break;
-            case ast::UnaryOp::kIndirection:
-                inst = builder_.Indirection(ty, val.Get());
-                break;
             case ast::UnaryOp::kNegation:
                 inst = builder_.Negation(ty, val.Get());
                 break;
diff --git a/src/tint/ir/from_program_binary_test.cc b/src/tint/ir/from_program_binary_test.cc
index a9b70e7e39..76e24fb10f 100644
--- a/src/tint/ir/from_program_binary_test.cc
+++ b/src/tint/ir/from_program_binary_test.cc
@@ -58,14 +58,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Increment) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = add %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:ptr<private, u32, read_write> = add %v1:ptr<private, u32, read_write>, 1u
+    store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -81,14 +81,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAdd) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = add %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:ptr<private, u32, read_write> = add %v1:ptr<private, u32, read_write>, 1u
+    store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -127,14 +127,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Decrement) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, i32, read_write> = var
+  %v1:ptr<private, i32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, i32, read_write> = sub %v1:ref<private, i32, read_write>, 1i
-    store %v1:ref<private, i32, read_write>, %2:ref<private, i32, read_write>
+    %2:ptr<private, i32, read_write> = sub %v1:ptr<private, i32, read_write>, 1u
+    store %v1:ptr<private, i32, read_write>, %2:ptr<private, i32, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -150,14 +150,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundSubtract) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = sub %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:ptr<private, u32, read_write> = sub %v1:ptr<private, u32, read_write>, 1u
+    store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -196,14 +196,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundMultiply) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = mul %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:ptr<private, u32, read_write> = mul %v1:ptr<private, u32, read_write>, 1u
+    store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -242,14 +242,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundDiv) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = div %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:ptr<private, u32, read_write> = div %v1:ptr<private, u32, read_write>, 1u
+    store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -288,14 +288,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundModulo) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = mod %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:ptr<private, u32, read_write> = mod %v1:ptr<private, u32, read_write>, 1u
+    store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -334,14 +334,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAnd) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, bool, read_write> = var
+  %v1:ptr<private, bool, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, bool, read_write> = and %v1:ref<private, bool, read_write>, false
-    store %v1:ref<private, bool, read_write>, %2:ref<private, bool, read_write>
+    %2:ptr<private, bool, read_write> = and %v1:ptr<private, bool, read_write>, false
+    store %v1:ptr<private, bool, read_write>, %2:ptr<private, bool, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -380,14 +380,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundOr) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, bool, read_write> = var
+  %v1:ptr<private, bool, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, bool, read_write> = or %v1:ref<private, bool, read_write>, false
-    store %v1:ref<private, bool, read_write>, %2:ref<private, bool, read_write>
+    %2:ptr<private, bool, read_write> = or %v1:ptr<private, bool, read_write>, false
+    store %v1:ptr<private, bool, read_write>, %2:ptr<private, bool, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -426,14 +426,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundXor) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = xor %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:ptr<private, u32, read_write> = xor %v1:ptr<private, u32, read_write>, 1u
+    store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -706,14 +706,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftLeft) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = shiftl %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:ptr<private, u32, read_write> = shiftl %v1:ptr<private, u32, read_write>, 1u
+    store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -752,14 +752,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftRight) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, u32, read_write> = var
+  %v1:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %2:ref<private, u32, read_write> = shiftr %v1:ref<private, u32, read_write>, 1u
-    store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+    %2:ptr<private, u32, read_write> = shiftr %v1:ptr<private, u32, read_write>, 1u
+    store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
   } -> %func_end # return
 } %func_end
 
diff --git a/src/tint/ir/from_program_call_test.cc b/src/tint/ir/from_program_call_test.cc
index 18fe3b1e70..f4c0ee79fc 100644
--- a/src/tint/ir/from_program_call_test.cc
+++ b/src/tint/ir/from_program_call_test.cc
@@ -100,13 +100,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %i:ref<private, i32, read_write> = var, 1i
+  %i:ptr<private, i32, read_write> = var, 1i
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %tint_symbol:f32 = convert i32, %i:ref<private, i32, read_write>
+    %tint_symbol:f32 = convert i32, %i:ptr<private, i32, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -121,7 +121,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %i:ref<private, vec3<f32>, read_write> = var, vec3<f32> 0.0f
+  %i:ptr<private, vec3<f32>, read_write> = var, vec3<f32> 0.0f
 }
 
 
@@ -137,13 +137,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %i:ref<private, f32, read_write> = var, 1.0f
+  %i:ptr<private, f32, read_write> = var, 1.0f
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %i:ref<private, f32, read_write>
+    %tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %i:ptr<private, f32, read_write>
   } -> %func_end # return
 } %func_end
 
diff --git a/src/tint/ir/from_program_store_test.cc b/src/tint/ir/from_program_store_test.cc
index 58fbe4aa5a..808138c424 100644
--- a/src/tint/ir/from_program_store_test.cc
+++ b/src/tint/ir/from_program_store_test.cc
@@ -36,13 +36,13 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Assign) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %a:ref<private, u32, read_write> = var
+  %a:ptr<private, u32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    store %a:ref<private, u32, read_write>, 4u
+    store %a:ptr<private, u32, read_write>, 4u
   } -> %func_end # return
 } %func_end
 
diff --git a/src/tint/ir/from_program_unary_test.cc b/src/tint/ir/from_program_unary_test.cc
index 3bcc5b1273..e386554235 100644
--- a/src/tint/ir/from_program_unary_test.cc
+++ b/src/tint/ir/from_program_unary_test.cc
@@ -105,13 +105,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_AddressOf) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, i32, read_write> = var
+  %v2:ptr<private, i32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %v2:ptr<private, i32, read_write> = addr_of %v1:ref<private, i32, read_write>
   } -> %func_end # return
 } %func_end
 
@@ -122,7 +121,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) {
     GlobalVar("v1", builtin::AddressSpace::kPrivate, ty.i32());
     utils::Vector stmts = {
         Decl(Let("v3", AddressOf("v1"))),
-        Decl(Let("v2", Deref("v3"))),
+        Assign(Deref("v3"), 42_i),
     };
     WrapInFunction(stmts);
 
@@ -130,14 +129,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %v1:ref<private, i32, read_write> = var
+  %v3:ptr<private, i32, read_write> = var
 }
 
 
 %fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn3 = block {
-    %v3:ptr<private, i32, read_write> = addr_of %v1:ref<private, i32, read_write>
-    %v2:i32 = indirection %v3:ptr<private, i32, read_write>
+    store %v3:ptr<private, i32, read_write>, 42i
   } -> %func_end # return
 } %func_end
 
diff --git a/src/tint/ir/from_program_var_test.cc b/src/tint/ir/from_program_var_test.cc
index 0cc9cbb894..e235f8880e 100644
--- a/src/tint/ir/from_program_var_test.cc
+++ b/src/tint/ir/from_program_var_test.cc
@@ -33,7 +33,7 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %a:ref<private, u32, read_write> = var
+  %a:ptr<private, u32, read_write> = var
 }
 
 
@@ -48,7 +48,7 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
-  %a:ref<private, u32, read_write> = var, 2u
+  %a:ptr<private, u32, read_write> = var, 2u
 }
 
 
@@ -65,7 +65,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn2 = block {
-    %a:ref<function, u32, read_write> = var
+    %a:ptr<function, u32, read_write> = var
   } -> %func_end # return
 } %func_end
 
@@ -83,7 +83,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn2 = block {
-    %a:ref<function, u32, read_write> = var, 2u
+    %a:ptr<function, u32, read_write> = var, 2u
   } -> %func_end # return
 } %func_end
 
diff --git a/src/tint/ir/to_program.cc b/src/tint/ir/to_program.cc
index 4d38015672..219490982a 100644
--- a/src/tint/ir/to_program.cc
+++ b/src/tint/ir/to_program.cc
@@ -272,12 +272,12 @@ class State {
 
     const ast::VariableDeclStatement* Var(const ir::Var* var) {
         Symbol name = NameOf(var);
-        auto* ptr = var->Type()->As<type::Reference>();
+        auto* ptr = var->Type()->As<type::Pointer>();
         if (!ptr) {
             Err("Incorrect type for var");
             return nullptr;
         }
-        auto ty = Type(ptr);
+        auto ty = Type(ptr->StoreType());
         const ast::Expression* init = nullptr;
         if (var->initializer) {
             init = Expr(var->initializer);
@@ -432,7 +432,10 @@ class State {
                                   : builtin::Access::kUndefined;
                 return b.ty.pointer(el.Get(), address_space, access);
             },
-            [&](const type::Reference* r) { return Type(r->StoreType()); },
+            [&](const type::Reference*) -> utils::Result<ast::Type> {
+                TINT_ICE(IR, b.Diagnostics()) << "reference types should never appear in the IR";
+                return ast::Type{};
+            },
             [&](Default) {
                 UNHANDLED_CASE(ty);
                 return ast::Type{};
diff --git a/src/tint/ir/unary.h b/src/tint/ir/unary.h
index e665386597..46edd450e9 100644
--- a/src/tint/ir/unary.h
+++ b/src/tint/ir/unary.h
@@ -25,9 +25,7 @@ class Unary : public utils::Castable<Unary, Instruction> {
   public:
     /// The kind of instruction.
     enum class Kind {
-        kAddressOf,
         kComplement,
-        kIndirection,
         kNegation,
     };
 
@@ -50,7 +48,7 @@ class Unary : public utils::Castable<Unary, Instruction> {
     const Value* Val() const { return val_; }
 
     /// the kind of unary instruction
-    Kind kind = Kind::kAddressOf;
+    Kind kind = Kind::kNegation;
 
     /// the result type of the instruction
     const type::Type* result_type = nullptr;
diff --git a/src/tint/ir/unary_test.cc b/src/tint/ir/unary_test.cc
index 392a75e340..280c7de831 100644
--- a/src/tint/ir/unary_test.cc
+++ b/src/tint/ir/unary_test.cc
@@ -23,27 +23,6 @@ using namespace tint::number_suffixes;  // NOLINT
 
 using IR_InstructionTest = TestHelper;
 
-TEST_F(IR_InstructionTest, CreateAddressOf) {
-    Module mod;
-    Builder b{mod};
-
-    // TODO(dsinclair): This would be better as an identifier, but works for now.
-    const auto* inst = b.AddressOf(
-        b.ir.types.Get<type::Pointer>(b.ir.types.Get<type::I32>(), builtin::AddressSpace::kPrivate,
-                                      builtin::Access::kReadWrite),
-        b.Constant(4_i));
-
-    ASSERT_TRUE(inst->Is<Unary>());
-    EXPECT_EQ(inst->kind, Unary::Kind::kAddressOf);
-
-    ASSERT_NE(inst->Type(), nullptr);
-
-    ASSERT_TRUE(inst->Val()->Is<Constant>());
-    auto lhs = inst->Val()->As<Constant>()->value;
-    ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
-    EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
-}
-
 TEST_F(IR_InstructionTest, CreateComplement) {
     Module mod;
     Builder b{mod};
@@ -58,22 +37,6 @@ TEST_F(IR_InstructionTest, CreateComplement) {
     EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
 }
 
-TEST_F(IR_InstructionTest, CreateIndirection) {
-    Module mod;
-    Builder b{mod};
-
-    // TODO(dsinclair): This would be better as an identifier, but works for now.
-    const auto* inst = b.Indirection(b.ir.types.Get<type::I32>(), b.Constant(4_i));
-
-    ASSERT_TRUE(inst->Is<Unary>());
-    EXPECT_EQ(inst->kind, Unary::Kind::kIndirection);
-
-    ASSERT_TRUE(inst->Val()->Is<Constant>());
-    auto lhs = inst->Val()->As<Constant>()->value;
-    ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
-    EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
-}
-
 TEST_F(IR_InstructionTest, CreateNegation) {
     Module mod;
     Builder b{mod};