[spirv-reader][ir] Support atomic access through a `let`.
Add support for accessing an atomic through a `let` instruction.
Bug: 404501988
Change-Id: Icbcf144f891a715e3a2e2fc942a7661bbf177a97
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/233074
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/reader/lower/atomics.cc b/src/tint/lang/spirv/reader/lower/atomics.cc
index ac91882..720b4e0 100644
--- a/src/tint/lang/spirv/reader/lower/atomics.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics.cc
@@ -124,6 +124,7 @@
tint::Switch( //
res->Instruction(), //
[&](core::ir::Access* a) { values_to_convert_.Push(a->Object()); }, //
+ [&](core::ir::Let* l) { values_to_convert_.Push(l->Value()); }, //
[&](core::ir::Var*) {}, //
TINT_ICE_ON_NO_MATCH);
}
diff --git a/src/tint/lang/spirv/reader/lower/atomics_test.cc b/src/tint/lang/spirv/reader/lower/atomics_test.cc
index d3bef8e..0c9c90d 100644
--- a/src/tint/lang/spirv/reader/lower/atomics_test.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics_test.cc
@@ -84,7 +84,7 @@
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_AtomicsTest, DISABLED_ArrayStore_CopiedObject) {
+TEST_F(SpirvReader_AtomicsTest, ArrayStore_CopiedObject) {
auto* f = b.ComputeFunction("main");
core::ir::Var* wg = nullptr;
@@ -117,7 +117,18 @@
Run(Atomics);
auto* expect = R"(
-UNIMPLEMENTED
+$B1: { # root
+ %wg:ptr<workgroup, array<atomic<u32>, 4>, read_write> = var undef
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %3:ptr<workgroup, array<atomic<u32>, 4>, read_write> = let %wg
+ %4:ptr<workgroup, atomic<u32>, read_write> = access %3, 1i
+ %5:void = atomicStore %4, 2u
+ ret
+ }
+}
)";
ASSERT_EQ(expect, str());
}
