[tint][ir] Fix indexing of abstract typed constants
Fixed: tint:1972
Change-Id: I33c60ce65de690ae306bcb3f71f464d3f2e5da32
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/139641
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index 69b80ea..19a16ea 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -867,6 +867,9 @@
std::vector<const ast::Expression*> accessors;
const ast::Expression* object = expr;
while (true) {
+ if (program_->Sem().GetVal(object)->ConstantValue()) {
+ break; // Reached a constant expression. Stop traversal.
+ }
if (auto* array = object->As<ast::IndexAccessorExpression>()) {
accessors.push_back(object);
object = array->object;
@@ -944,7 +947,7 @@
}
bool GenerateMemberAccessor(const ast::MemberAccessorExpression* expr, AccessorInfo& info) {
- auto* expr_sem = program_->Sem().Get(expr)->UnwrapLoad();
+ auto* expr_sem = program_->Sem().Get(expr)->Unwrap();
return tint::Switch(
expr_sem, //
@@ -966,7 +969,7 @@
// intermediate steps need different result types.
auto* result_type = info.result_type;
- // Emit any preceeding member/index accessors
+ // Emit any preceding member/index accessors
if (!info.indices.IsEmpty()) {
// The access chain is being split, the initial part of than will have a
// resulting type that matches the object being swizzled.
diff --git a/src/tint/ir/from_program_accessor_test.cc b/src/tint/ir/from_program_accessor_test.cc
index 3628f10..cd470c0 100644
--- a/src/tint/ir/from_program_accessor_test.cc
+++ b/src/tint/ir/from_program_accessor_test.cc
@@ -35,7 +35,7 @@
auto* a = Var("a", ty.vec3<u32>(), builtin::AddressSpace::kFunction);
auto* expr = Decl(Let("b", IndexAccessor(a, 2_u)));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -58,7 +58,7 @@
auto* a = Var("a", ty.mat3x4<f32>(), builtin::AddressSpace::kFunction);
auto* expr = Decl(Let("b", IndexAccessor(IndexAccessor(a, 2_u), 3_u)));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -85,7 +85,7 @@
});
auto* a = Var("a", ty.Of(s), builtin::AddressSpace::kFunction);
auto* expr = Decl(Let("b", MemberAccessor(a, "foo")));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -121,7 +121,7 @@
});
auto* a = Var("a", ty.Of(outer), builtin::AddressSpace::kFunction);
auto* expr = Decl(Let("b", MemberAccessor(MemberAccessor(a, "foo"), "bar")));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -166,7 +166,7 @@
auto* expr = Decl(Let(
"b",
MemberAccessor(IndexAccessor(MemberAccessor(IndexAccessor(a, 0_u), "foo"), 1_u), "bar")));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -200,7 +200,7 @@
auto* a = Var("a", ty.array<u32, 4>(), builtin::AddressSpace::kFunction);
auto* assign = Assign(IndexAccessor(a, 2_u), 0_u);
- WrapInFunction(Block(utils::Vector{Decl(a), assign}));
+ WrapInFunction(Decl(a), assign);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -223,7 +223,7 @@
auto* a = Var("a", ty.vec2<f32>(), builtin::AddressSpace::kFunction);
auto* expr = Decl(Let("b", MemberAccessor(a, "y")));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -246,7 +246,7 @@
auto* a = Var("a", ty.vec3<f32>(), builtin::AddressSpace::kFunction);
auto* expr = Decl(Let("b", MemberAccessor(a, "zyxz")));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -269,7 +269,7 @@
auto* a = Var("a", ty.vec3<f32>(), builtin::AddressSpace::kFunction);
auto* expr = Decl(Let("b", MemberAccessor(MemberAccessor(a, "zyx"), "yy")));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -300,7 +300,7 @@
auto* expr = Decl(Let(
"b",
IndexAccessor(MemberAccessor(MemberAccessor(MemberAccessor(a, "foo"), "zyx"), "yx"), 0_u)));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -330,7 +330,7 @@
// let b = a[2]
auto* a = Let("a", ty.vec3<u32>(), vec(ty.u32(), 3));
auto* expr = Decl(Let("b", IndexAccessor(a, 2_u)));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -351,7 +351,7 @@
auto* a = Let("a", ty.mat3x4<f32>(), Call<mat3x4<f32>>());
auto* expr = Decl(Let("b", IndexAccessor(IndexAccessor(a, 2_u), 3_u)));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -376,7 +376,7 @@
});
auto* a = Let("a", ty.Of(s), Call("MyStruct"));
auto* expr = Decl(Let("b", MemberAccessor(a, "foo")));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -410,7 +410,7 @@
});
auto* a = Let("a", ty.Of(outer), Call("Outer"));
auto* expr = Decl(Let("b", MemberAccessor(MemberAccessor(a, "foo"), "bar")));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -453,7 +453,7 @@
auto* expr = Decl(Let(
"b",
MemberAccessor(IndexAccessor(MemberAccessor(IndexAccessor(a, 0_u), "foo"), 1_u), "bar")));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -485,7 +485,7 @@
auto* a = Let("a", ty.vec2<f32>(), vec(ty.f32(), 2));
auto* expr = Decl(Let("b", MemberAccessor(a, "y")));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -506,7 +506,7 @@
auto* a = Let("a", ty.vec3<f32>(), vec(ty.f32(), 3));
auto* expr = Decl(Let("b", MemberAccessor(a, "zyxz")));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -527,7 +527,7 @@
auto* a = Let("a", ty.vec3<f32>(), vec(ty.f32(), 3));
auto* expr = Decl(Let("b", MemberAccessor(MemberAccessor(a, "zyx"), "yy")));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -556,7 +556,7 @@
auto* expr = Decl(Let(
"b",
IndexAccessor(MemberAccessor(MemberAccessor(MemberAccessor(a, "foo"), "zyx"), "yx"), 0_u)));
- WrapInFunction(Block(utils::Vector{Decl(a), expr}));
+ WrapInFunction(Decl(a), expr);
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
@@ -579,5 +579,53 @@
)");
}
+TEST_F(IR_FromProgramAccessorTest, Accessor_Const_AbstractVectorWithIndex) {
+ // const v = vec3(1, 2, 3);
+ // let i = 1;
+ // var b = v[i];
+
+ auto* v = Const("v", Call<vec3<Infer>>(1_a, 2_a, 3_a));
+ auto* i = Let("i", Expr(1_i));
+ auto* b = Var("b", IndexAccessor("v", "i"));
+ WrapInFunction(v, i, b);
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %2:i32 = access vec3<i32>(1i, 2i, 3i), 1i
+ %b:ptr<function, i32, read_write> = var, %2
+ ret
+ }
+}
+)");
+}
+
+TEST_F(IR_FromProgramAccessorTest, Accessor_Const_AbstractVectorWithSwizzleAndIndex) {
+ // const v = vec3(1, 2, 3);
+ // let i = 1;
+ // var b = v.rg[i];
+
+ auto* v = Const("v", Call<vec3<Infer>>(1_a, 2_a, 3_a));
+ auto* i = Let("i", Expr(1_i));
+ auto* b = Var("b", IndexAccessor(MemberAccessor("v", "rg"), "i"));
+ WrapInFunction(v, i, b);
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %2:i32 = access vec2<i32>(1i, 2i), 1i
+ %b:ptr<function, i32, read_write> = var, %2
+ ret
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::ir
