diff --git a/src/tint/ir/builder_impl.cc b/src/tint/ir/builder_impl.cc
index ecf7a30029..bf6a0e725f 100644
--- a/src/tint/ir/builder_impl.cc
+++ b/src/tint/ir/builder_impl.cc
@@ -775,9 +775,7 @@ void BuilderImpl::EmitVariable(const ast::Variable* var) {
                 if (!init) {
                     return;
                 }
-
-                auto* store = builder.Store(val, init.Get());
-                current_flow_block->instructions.Push(store);
+                val->initializer = init.Get();
             }
             // Store the declaration so we can get the instruction to store too
             scopes_.Set(v->name->symbol, val);
diff --git a/src/tint/ir/builder_impl_binary_test.cc b/src/tint/ir/builder_impl_binary_test.cc
index 20c5cca96f..3be8405e09 100644
--- a/src/tint/ir/builder_impl_binary_test.cc
+++ b/src/tint/ir/builder_impl_binary_test.cc
@@ -54,7 +54,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAdd) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%v1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
 
 
 
@@ -96,7 +96,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundSubtract) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%v1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
 
 
 
@@ -138,7 +138,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundMultiply) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%v1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
 
 
 
@@ -180,7 +180,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundDiv) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%v1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
 
 
 
@@ -222,7 +222,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundModulo) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%v1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
 
 
 
@@ -264,7 +264,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAnd) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%v1:ref<private, bool, read_write> = var private read_write
+%v1:ref<private, bool, read_write> = var private, read_write
 
 
 
@@ -306,7 +306,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundOr) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%v1:ref<private, bool, read_write> = var private read_write
+%v1:ref<private, bool, read_write> = var private, read_write
 
 
 
@@ -348,7 +348,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundXor) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%v1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
 
 
 
@@ -379,7 +379,7 @@ func_end
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
   %fn4 = block
   %1:bool = call my_func
-  %tint_symbol:bool = var function read_write
+  %tint_symbol:bool = var function, read_write
   store %tint_symbol:bool, %1:bool
   branch %fn5
 
@@ -414,7 +414,7 @@ func_end
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
   %fn4 = block
   %1:bool = call my_func
-  %tint_symbol:bool = var function read_write
+  %tint_symbol:bool = var function, read_write
   store %tint_symbol:bool, %1:bool
   branch %fn5
 
@@ -569,7 +569,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftLeft) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%v1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
 
 
 
@@ -611,7 +611,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftRight) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%v1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
 
 
 
@@ -645,7 +645,7 @@ func_end
   %fn4 = block
   %1:f32 = call my_func
   %2:bool = lt %1:f32, 2.0f
-  %tint_symbol:bool = var function read_write
+  %tint_symbol:bool = var function, read_write
   store %tint_symbol:bool, %2:bool
   branch %fn5
 
diff --git a/src/tint/ir/builder_impl_call_test.cc b/src/tint/ir/builder_impl_call_test.cc
index 2e63a212ec..0fb070e2e8 100644
--- a/src/tint/ir/builder_impl_call_test.cc
+++ b/src/tint/ir/builder_impl_call_test.cc
@@ -92,8 +92,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
     ASSERT_TRUE(r);
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%i:ref<private, i32, read_write> = var private read_write
-store %i:ref<private, i32, read_write>, 1i
+%i:ref<private, i32, read_write> = var private, read_write, 1i
 
 
 
@@ -116,8 +115,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) {
     ASSERT_TRUE(r);
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%i:ref<private, vec3<f32>, read_write> = var private read_write
-store %i:ref<private, vec3<f32>, read_write>, vec3<f32> 0.0f
+%i:ref<private, vec3<f32>, read_write> = var private, read_write, vec3<f32> 0.0f
 
 
 
@@ -135,8 +133,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
     ASSERT_TRUE(r);
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%i:ref<private, f32, read_write> = var private read_write
-store %i:ref<private, f32, read_write>, 1.0f
+%i:ref<private, f32, read_write> = var private, read_write, 1.0f
 
 
 
diff --git a/src/tint/ir/builder_impl_store_test.cc b/src/tint/ir/builder_impl_store_test.cc
index 75bb7c0416..82fbc0a45f 100644
--- a/src/tint/ir/builder_impl_store_test.cc
+++ b/src/tint/ir/builder_impl_store_test.cc
@@ -37,7 +37,7 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Assign) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%a:ref<private, u32, read_write> = var private read_write
+%a:ref<private, u32, read_write> = var private, read_write
 
 
 
diff --git a/src/tint/ir/builder_impl_unary_test.cc b/src/tint/ir/builder_impl_unary_test.cc
index 16ec6f9dda..2323acbe0d 100644
--- a/src/tint/ir/builder_impl_unary_test.cc
+++ b/src/tint/ir/builder_impl_unary_test.cc
@@ -91,7 +91,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_AddressOf) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%v1:ref<private, i32, read_write> = var private read_write
+%v1:ref<private, i32, read_write> = var private, read_write
 
 
 
@@ -117,7 +117,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%v1:ref<private, i32, read_write> = var private read_write
+%v1:ref<private, i32, read_write> = var private, read_write
 
 
 
diff --git a/src/tint/ir/builder_impl_var_test.cc b/src/tint/ir/builder_impl_var_test.cc
index 5d96a279e7..a3f4d27fee 100644
--- a/src/tint/ir/builder_impl_var_test.cc
+++ b/src/tint/ir/builder_impl_var_test.cc
@@ -34,7 +34,7 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%a:ref<private, u32, read_write> = var private read_write
+%a:ref<private, u32, read_write> = var private, read_write
 
 
 
@@ -50,8 +50,7 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) {
     auto m = r.Move();
 
     EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%a:ref<private, u32, read_write> = var private read_write
-store %a:ref<private, u32, read_write>, 2u
+%a:ref<private, u32, read_write> = var private, read_write, 2u
 
 
 
@@ -69,7 +68,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
     EXPECT_EQ(Disassemble(m),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
   %fn2 = block
-  %a:ref<function, u32, read_write> = var function read_write
+  %a:ref<function, u32, read_write> = var function, read_write
   ret
 func_end
 
@@ -88,8 +87,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
     EXPECT_EQ(Disassemble(m),
               R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
   %fn2 = block
-  %a:ref<function, u32, read_write> = var function read_write
-  store %a:ref<function, u32, read_write>, 2u
+  %a:ref<function, u32, read_write> = var function, read_write, 2u
   ret
 func_end
 
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 2d8d401fc4..600f5870b4 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -406,7 +406,11 @@ void Disassembler::EmitInstruction(const Instruction* inst) {
         },
         [&](const ir::Var* v) {
             EmitValue(v);
-            out_ << " = var " << v->address_space << " " << v->access;
+            out_ << " = var " << v->address_space << ", " << v->access;
+            if (v->initializer) {
+                out_ << ", ";
+                EmitValue(v->initializer);
+            }
         });
 }
 
diff --git a/src/tint/ir/var.h b/src/tint/ir/var.h
index 1387c609d4..8ac39d2c35 100644
--- a/src/tint/ir/var.h
+++ b/src/tint/ir/var.h
@@ -42,6 +42,9 @@ class Var : public utils::Castable<Var, Instruction> {
 
     /// The variable access mode
     builtin::Access access = builtin::Access::kUndefined;
+
+    /// The optional initializer
+    Value* initializer = nullptr;
 };
 
 }  // namespace tint::ir