commit cde3c148dd0f34f8819577b27ac8ee6b10be751d
author dan sinclair <dsinclair@chromium.org> Mon Apr 14 15:51:33 2025 -0700
committer Dawn LUCI CQ <dawn-scoped@luci-project-accounts.iam.gserviceaccount.com> Mon Apr 14 15:51:33 2025 -0700
tree 32d0a7d556806e5ee62cbbb30434f30736969fce
parent 5d3c16fde743b7517b9209b6ebf85fe33d2d534f

[spirv-reader][ir] Enable atomic store test.

Enable disabled atomicStore test which passes correctly.

Bug: 391487430
Change-Id: Idcb6566cfacb695de08bc69d70afb8b171076d22
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/236914
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>

src/tint/lang/spirv/reader/lower/atomics_test.cc

diff --git a/src/tint/lang/spirv/reader/lower/atomics_test.cc b/src/tint/lang/spirv/reader/lower/atomics_test.cc
index ed4615a..1de66a5 100644
--- a/src/tint/lang/spirv/reader/lower/atomics_test.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics_test.cc
@@ -1629,7 +1629,7 @@
     ASSERT_EQ(expect, str());
 }
 
-TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicStore) {
+TEST_F(SpirvReader_AtomicsTest, AtomicStore) {
     auto* f = b.ComputeFunction("main");
 
     auto* sb = ty.Struct(mod.symbols.New("S"), {
@@ -1655,11 +1655,11 @@
 
         auto* a1 = b.Access(ty.ptr<storage, u32, read_write>(), sg, 1_u);
         b.Call<spirv::ir::BuiltinCall>(ty.void_(), spirv::BuiltinFn::kAtomicStore, a1, 1_u, 0_u,
-                                       1_u);
+                                       2_u);
         b.Call<spirv::ir::BuiltinCall>(ty.void_(), spirv::BuiltinFn::kAtomicStore, wg_i32, 1_u, 0_u,
-                                       1_i);
+                                       3_i);
         b.Call<spirv::ir::BuiltinCall>(ty.void_(), spirv::BuiltinFn::kAtomicStore, wg_u32, 1_u, 0_u,
-                                       1_u);
+                                       4_u);
         b.Return(f);
     });
 
@@ -1680,9 +1680,9 @@
     %5:ptr<storage, i32, read_write> = access %sb, 0u
     %6:void = spirv.atomic_store %5, 1u, 0u, 1i
     %7:ptr<storage, u32, read_write> = access %sb, 1u
-    %8:void = spirv.atomic_store %7, 1u, 0u, 1u
-    %9:void = spirv.atomic_store %wg_i32, 1u, 0u, 1i
-    %10:void = spirv.atomic_store %wg_u32, 1u, 0u, 1u
+    %8:void = spirv.atomic_store %7, 1u, 0u, 2u
+    %9:void = spirv.atomic_store %wg_i32, 1u, 0u, 3i
+    %10:void = spirv.atomic_store %wg_u32, 1u, 0u, 4u
     ret
   }
 }
@@ -1692,7 +1692,33 @@
     Run(Atomics);
 
     auto* expect = R"(
-UNIMPLEMENTED
+S = struct @align(4) {
+  a:i32 @offset(0)
+  b:u32 @offset(4)
+}
+
+S_atomic = struct @align(4) {
+  a:atomic<i32> @offset(0)
+  b:atomic<u32> @offset(4)
+}
+
+$B1: {  # root
+  %sb:ptr<storage, S_atomic, read_write> = var undef @binding_point(0, 0)
+  %wg_i32:ptr<workgroup, atomic<i32>, read_write> = var undef
+  %wg_u32:ptr<workgroup, atomic<u32>, read_write> = var undef
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B2: {
+    %5:ptr<storage, atomic<i32>, read_write> = access %sb, 0u
+    %6:void = atomicStore %5, 1i
+    %7:ptr<storage, atomic<u32>, read_write> = access %sb, 1u
+    %8:void = atomicStore %7, 2u
+    %9:void = atomicStore %wg_i32, 3i
+    %10:void = atomicStore %wg_u32, 4u
+    ret
+  }
+}
 )";
     ASSERT_EQ(expect, str());
 }