diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc index 4773f3fcbb..20555d0e5e 100644 --- a/src/tint/ir/from_program.cc +++ b/src/tint/ir/from_program.cc @@ -842,8 +842,6 @@ class Impl { ref->Access()); auto* val = builder_.Declare(ty); - current_flow_block_->Instructions().Push(val); - if (v->initializer) { auto init = EmitExpression(v->initializer); if (!init) { @@ -851,6 +849,8 @@ class Impl { } val->SetInitializer(init.Get()); } + current_flow_block_->Instructions().Push(val); + // Store the declaration so we can get the instruction to store too scopes_.Set(v->name->symbol, val); diff --git a/src/tint/ir/from_program_var_test.cc b/src/tint/ir/from_program_var_test.cc index 6cecf47d1b..c533738af2 100644 --- a/src/tint/ir/from_program_var_test.cc +++ b/src/tint/ir/from_program_var_test.cc @@ -72,7 +72,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) { )"); } -TEST_F(IR_BuilderImplTest, Emit_Var_Init) { +TEST_F(IR_BuilderImplTest, Emit_Var_Init_Constant) { auto* expr = Expr(2_u); auto* a = Var("a", ty.u32(), builtin::AddressSpace::kFunction, expr); WrapInFunction(a); @@ -89,5 +89,27 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) { } )"); } + +TEST_F(IR_BuilderImplTest, Emit_Var_Init_NonConstant) { + auto* a = Var("a", ty.u32(), builtin::AddressSpace::kFunction); + auto* b = Var("b", ty.u32(), builtin::AddressSpace::kFunction, Add("a", 2_u)); + WrapInFunction(a, b); + + auto m = Build(); + ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + + EXPECT_EQ(Disassemble(m.Get()), + R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { + %b1 = block { + %a:ptr = var + %3:u32 = load %a + %4:u32 = add %3, 2u + %b:ptr = var, %4 + ret + } +} +)"); +} + } // namespace } // namespace tint::ir