[ir] Emit initializer before var declaration

Otherwise for non-constant initializers the var instruction will
reference the result of an instruction that comes after it.

Bug: tint:1718
Change-Id: Ie8c6c900768277c344f1e1304b0812a546c889a6
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/134560
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
This commit is contained in:
James Price 2023-05-26 13:45:10 +00:00 committed by Dawn LUCI CQ
parent 0202531159
commit 2ee63ffc0c
2 changed files with 25 additions and 3 deletions

View File

@ -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);

View File

@ -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<function, u32, read_write> = var
%3:u32 = load %a
%4:u32 = add %3, 2u
%b:ptr<function, u32, read_write> = var, %4
ret
}
}
)");
}
} // namespace
} // namespace tint::ir