[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 <[email protected]>
Commit-Queue: dan sinclair <[email protected]>

Author: dj2

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

@@ -1629,7 +1629,7 @@ UNIMPLEMENTED
     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 @@ TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicStore) {
 
         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 @@ S = struct @align(4) {
     %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 @@ S = struct @align(4) {
     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());
 }