[spirv-reader] Avoid creating access with no indices
If an `OpAccessChain` instruction has no indices, just forward the
source object instead of creating an empty `access` instruction.
Simplify the two downstream transforms accordingly.
Bug: 370743660
Change-Id: I26e9c5c09d4b22cdeb45c6d2f08727fce0c46f37
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/211097
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/lower/shader_io.cc b/src/tint/lang/spirv/reader/lower/shader_io.cc
index 7fc8d18..24d4601 100644
--- a/src/tint/lang/spirv/reader/lower/shader_io.cc
+++ b/src/tint/lang/spirv/reader/lower/shader_io.cc
@@ -280,16 +280,10 @@
to_destroy.Push(lve);
},
[&](core::ir::Access* a) {
- if (!a->Indices().IsEmpty()) {
- // Remove the pointer from the source and destination type.
- a->SetOperand(core::ir::Access::kObjectOperandOffset, object);
- a->Result(0)->SetType(a->Result(0)->Type()->UnwrapPtr());
- ReplaceInputPointerUses(var, a->Result(0));
- } else {
- // Fold the access away and replace its uses.
- ReplaceInputPointerUses(var, a->Result(0));
- to_destroy.Push(a);
- }
+ // Remove the pointer from the source and destination type.
+ a->SetOperand(core::ir::Access::kObjectOperandOffset, object);
+ a->Result(0)->SetType(a->Result(0)->Type()->UnwrapPtr());
+ ReplaceInputPointerUses(var, a->Result(0));
},
TINT_ICE_ON_NO_MATCH);
});
diff --git a/src/tint/lang/spirv/reader/lower/shader_io_test.cc b/src/tint/lang/spirv/reader/lower/shader_io_test.cc
index 8afe5e2..9647fbf 100644
--- a/src/tint/lang/spirv/reader/lower/shader_io_test.cc
+++ b/src/tint/lang/spirv/reader/lower/shader_io_test.cc
@@ -645,59 +645,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvReader_ShaderIOTest, Input_AccessChains) {
- auto* lid = b.Var("lid", ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()));
- {
- core::IOAttributes attributes;
- attributes.builtin = core::BuiltinValue::kLocalInvocationId;
- lid->SetAttributes(std::move(attributes));
- }
- mod.root_block->Append(lid);
-
- auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute);
- ep->SetWorkgroupSize(1, 1, 1);
- b.Append(ep->Block(), [&] {
- auto* access_1 = b.Access(ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()), lid);
- auto* access_2 = b.Access(ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()), access_1);
- auto* vec = b.Load(access_2);
- auto* z = b.LoadVectorElement(access_2, 2_u);
- b.Multiply<vec3<u32>>(vec, z);
- b.Return(ep);
- });
-
- auto* src = R"(
-$B1: { # root
- %lid:ptr<__in, vec3<u32>, read> = var @builtin(local_invocation_id)
-}
-
-%foo = @compute @workgroup_size(1, 1, 1) func():void {
- $B2: {
- %3:ptr<__in, vec3<u32>, read> = access %lid
- %4:ptr<__in, vec3<u32>, read> = access %3
- %5:vec3<u32> = load %4
- %6:u32 = load_vector_element %4, 2u
- %7:vec3<u32> = mul %5, %6
- ret
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-%foo = @compute @workgroup_size(1, 1, 1) func(%lid:vec3<u32> [@local_invocation_id]):void {
- $B1: {
- %3:u32 = access %lid, 2u
- %4:vec3<u32> = mul %lid, %3
- ret
- }
-}
-)";
-
- Run(ShaderIO);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(SpirvReader_ShaderIOTest, Inputs_Struct_LocationOnEachMember) {
auto* colors_str = ty.Struct(
mod.symbols.New("Colors"),
@@ -2035,72 +1982,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvReader_ShaderIOTest, Output_AccessChain) {
- auto* color = b.Var("color", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
- {
- core::IOAttributes attributes;
- attributes.location = 1u;
- color->SetAttributes(std::move(attributes));
- }
- mod.root_block->Append(color);
-
- auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
- b.Append(ep->Block(), [&] { //
- auto* access_1 = b.Access(ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()), color);
- auto* access_2 = b.Access(ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()), access_1);
- auto* load = b.LoadVectorElement(access_2, 2_u);
- auto* mul = b.Multiply<vec4<f32>>(b.Splat<vec4<f32>>(1_f), load);
- b.Store(access_2, mul);
- b.Return(ep);
- });
-
- auto* src = R"(
-$B1: { # root
- %color:ptr<__out, vec4<f32>, read_write> = var @location(1)
-}
-
-%foo = @fragment func():void {
- $B2: {
- %3:ptr<__out, vec4<f32>, read_write> = access %color
- %4:ptr<__out, vec4<f32>, read_write> = access %3
- %5:f32 = load_vector_element %4, 2u
- %6:vec4<f32> = mul vec4<f32>(1.0f), %5
- store %4, %6
- ret
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-$B1: { # root
- %color:ptr<private, vec4<f32>, read_write> = var
-}
-
-%foo_inner = func():void {
- $B2: {
- %3:ptr<private, vec4<f32>, read_write> = access %color
- %4:ptr<private, vec4<f32>, read_write> = access %3
- %5:f32 = load_vector_element %4, 2u
- %6:vec4<f32> = mul vec4<f32>(1.0f), %5
- store %4, %6
- ret
- }
-}
-%foo = @fragment func():vec4<f32> [@location(1)] {
- $B3: {
- %8:void = call %foo_inner
- %9:vec4<f32> = load %color
- ret %9
- }
-}
-)";
-
- Run(ShaderIO);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(SpirvReader_ShaderIOTest, Inputs_And_Outputs) {
auto* position = b.Var("position", ty.ptr(core::AddressSpace::kIn, ty.vec4<f32>()));
{
diff --git a/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc b/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc
index 97acad5..fa84275 100644
--- a/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc
+++ b/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc
@@ -140,11 +140,6 @@
sve->InsertBefore(store);
to_destroy.Push(store);
},
- [&](core::ir::Access* noop_access) {
- TINT_ASSERT(noop_access->Indices().IsEmpty());
- ReplaceAccessUses(noop_access, object, index);
- to_destroy.Push(noop_access);
- },
TINT_ICE_ON_NO_MATCH);
});
diff --git a/src/tint/lang/spirv/reader/lower/vector_element_pointer_test.cc b/src/tint/lang/spirv/reader/lower/vector_element_pointer_test.cc
index 640563f..710560f 100644
--- a/src/tint/lang/spirv/reader/lower/vector_element_pointer_test.cc
+++ b/src/tint/lang/spirv/reader/lower/vector_element_pointer_test.cc
@@ -65,64 +65,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvReader_VectorElementPointerTest, Access_NoIndices) {
- auto* foo = b.Function("foo", ty.vec4<u32>());
- b.Append(foo->Block(), [&] {
- auto* vec = b.Var<function, vec4<u32>>("vec");
- auto* access = b.Access<ptr<function, vec4<u32>>>(vec);
- b.Return(foo, b.Load(access));
- });
-
- auto* src = R"(
-%foo = func():vec4<u32> {
- $B1: {
- %vec:ptr<function, vec4<u32>, read_write> = var
- %3:ptr<function, vec4<u32>, read_write> = access %vec
- %4:vec4<u32> = load %3
- ret %4
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = src;
-
- Run(VectorElementPointer);
-
- EXPECT_EQ(expect, str());
-}
-
-TEST_F(SpirvReader_VectorElementPointerTest, Access_NoIndices_Chain) {
- auto* foo = b.Function("foo", ty.vec4<u32>());
- b.Append(foo->Block(), [&] {
- auto* vec = b.Var<function, vec4<u32>>("vec");
- auto* access_1 = b.Access<ptr<function, vec4<u32>>>(vec);
- auto* access_2 = b.Access<ptr<function, vec4<u32>>>(access_1);
- auto* access_3 = b.Access<ptr<function, vec4<u32>>>(access_2);
- b.Return(foo, b.Load(access_3));
- });
-
- auto* src = R"(
-%foo = func():vec4<u32> {
- $B1: {
- %vec:ptr<function, vec4<u32>, read_write> = var
- %3:ptr<function, vec4<u32>, read_write> = access %vec
- %4:ptr<function, vec4<u32>, read_write> = access %3
- %5:ptr<function, vec4<u32>, read_write> = access %4
- %6:vec4<u32> = load %5
- ret %6
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = src;
-
- Run(VectorElementPointer);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(SpirvReader_VectorElementPointerTest, Access_Component_NoUse) {
auto* foo = b.Function("foo", ty.void_());
b.Append(foo->Block(), [&] {
@@ -278,44 +220,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvReader_VectorElementPointerTest, AccessBeforeUse) {
- auto* foo = b.Function("foo", ty.void_());
- b.Append(foo->Block(), [&] {
- auto* vec = b.Var<function, vec4<u32>>("vec");
- auto* access_1 = b.Access<ptr<function, u32>>(vec, 2_u);
- auto* access_2 = b.Access<ptr<function, u32>>(access_1);
- b.Store(access_2, 42_u);
- b.Return(foo);
- });
-
- auto* src = R"(
-%foo = func():void {
- $B1: {
- %vec:ptr<function, vec4<u32>, read_write> = var
- %3:ptr<function, u32, read_write> = access %vec, 2u
- %4:ptr<function, u32, read_write> = access %3
- store %4, 42u
- ret
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-%foo = func():void {
- $B1: {
- %vec:ptr<function, vec4<u32>, read_write> = var
- store_vector_element %vec, 2u, 42u
- ret
- }
-}
-)";
-
- Run(VectorElementPointer);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(SpirvReader_VectorElementPointerTest, MultipleUses) {
auto* foo = b.Function("foo", ty.void_());
b.Append(foo->Block(), [&] {
diff --git a/src/tint/lang/spirv/reader/parser/memory_test.cc b/src/tint/lang/spirv/reader/parser/memory_test.cc
index 90e4945..9133ce5 100644
--- a/src/tint/lang/spirv/reader/parser/memory_test.cc
+++ b/src/tint/lang/spirv/reader/parser/memory_test.cc
@@ -726,9 +726,7 @@
%main = @compute @workgroup_size(1, 1, 1) func():void {
$B1: {
%2:ptr<function, u32, read_write> = var
- %3:ptr<function, u32, read_write> = access %2
- %4:ptr<function, u32, read_write> = access %3
- %5:u32 = load %4
+ %3:u32 = load %2
ret
}
)");
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 9a4094f..81117c8 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -549,6 +549,12 @@
}
auto* base = Value(inst.GetSingleWordOperand(2));
+ if (indices.IsEmpty()) {
+ // There are no indices, so just forward the base object.
+ AddValue(inst.result_id(), base);
+ return;
+ }
+
// Propagate the access mode of the base object.
auto access_mode = core::Access::kUndefined;
if (auto* ptr = base->Type()->As<core::type::Pointer>()) {
