[spirv-reader] Convert access results for atomics.

When converting an `Access` chain which is being turned into an atomic,
the result of the access also needs to be updated as those uses of the
access chain need to be atomics.

Fixed: 455844215
Change-Id: Ic183ae7774b780b48d216c8ab8b041ff74f1d734
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/269374
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: James Price <jrprice@google.com>
Auto-Submit: dan sinclair <dsinclair@chromium.org>

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

diff --git a/src/tint/lang/spirv/reader/lower/atomics.cc b/src/tint/lang/spirv/reader/lower/atomics.cc
index 5f825ba..dfa14ae 100644
--- a/src/tint/lang/spirv/reader/lower/atomics.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics.cc
@@ -319,6 +319,7 @@
                     [&](core::ir::Access* a) {
                         CheckForStructForking(a);
                         values_to_convert_.Push(a->Object());
+                        values_to_fix_usages_.Push(a->Result());
                     },
                     [&](core::ir::Let* l) {
                         values_to_convert_.Push(l->Value());

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 50d1c76..0592c2d 100644
--- a/src/tint/lang/spirv/reader/lower/atomics_test.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics_test.cc
@@ -3255,5 +3255,80 @@
     ASSERT_EQ(expect, str());
 }
 
+TEST_F(SpirvReader_AtomicsTest, RewriteLoad_SameAccessChain) {
+    auto* f = b.ComputeFunction("main");
+    core::ir::Var* buffer = nullptr;
+    b.Append(mod.root_block,
+             [&] {  //
+                 buffer = b.Var("buffer", ty.ptr<storage, array<u32>, read_write>());
+                 buffer->SetBindingPoint(0u, 0u);
+             });
+
+    b.Append(f->Block(), [&] {  //
+        auto* a = b.Access(ty.ptr<storage, u32, read_write>(), buffer, 1_i);
+        b.Let("x", b.Load(a));
+
+        auto* if_ = b.If(true);
+        b.Append(if_->True(), [&] {
+            b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicCompareExchange, a,
+                                           1_u, 0_u, 0_u, 0_u, 0_u);
+            b.ExitIf(if_);
+        });
+
+        b.Return(f);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %buffer:ptr<storage, array<u32>, read_write> = var undef @binding_point(0, 0)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B2: {
+    %3:ptr<storage, u32, read_write> = access %buffer, 1i
+    %4:u32 = load %3
+    %x:u32 = let %4
+    if true [t: $B3] {  # if_1
+      $B3: {  # true
+        %6:u32 = spirv.atomic_compare_exchange %3, 1u, 0u, 0u, 0u, 0u
+        exit_if  # if_1
+      }
+    }
+    ret
+  }
+}
+)";
+    ASSERT_EQ(src, str());
+    Run(Atomics);
+
+    auto* expect = R"(
+__atomic_compare_exchange_result_u32 = struct @align(4) {
+  old_value:u32 @offset(0)
+  exchanged:bool @offset(4)
+}
+
+$B1: {  # root
+  %buffer:ptr<storage, array<atomic<u32>>, read_write> = var undef @binding_point(0, 0)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B2: {
+    %3:ptr<storage, atomic<u32>, read_write> = access %buffer, 1i
+    %4:u32 = atomicLoad %3
+    %x:u32 = let %4
+    if true [t: $B3] {  # if_1
+      $B3: {  # true
+        %6:__atomic_compare_exchange_result_u32 = atomicCompareExchangeWeak %3, 0u, 0u
+        %7:u32 = access %6, 0u
+        exit_if  # if_1
+      }
+    }
+    ret
+  }
+}
+)";
+    ASSERT_EQ(expect, str());
+}
+
 }  // namespace
 }  // namespace tint::spirv::reader::lower

test/tint/bug/tint/455844215.spvasm

diff --git a/test/tint/bug/tint/455844215.spvasm b/test/tint/bug/tint/455844215.spvasm
new file mode 100644
index 0000000..fa0dfe1
--- /dev/null
+++ b/test/tint/bug/tint/455844215.spvasm
@@ -0,0 +1,53 @@
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 64 1 1
+
+               OpName %buffer_type "buffer_type"
+               OpName %buffer "buffer"
+               OpName %main "main"
+               OpDecorate %buffer DescriptorSet 0
+               OpDecorate %buffer Binding 0
+               OpDecorate %runtime_arr ArrayStride 4
+               OpMemberDecorate %buffer_type 0 Offset 0
+               OpDecorate %buffer_type BufferBlock
+
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+%runtime_arr = OpTypeRuntimeArray %uint
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+       %bool = OpTypeBool
+%buffer_type = OpTypeStruct %runtime_arr
+
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+       %true = OpConstantTrue %bool
+
+%ptr_buffer_type = OpTypePointer Uniform %buffer_type
+   %ptr_uint = OpTypePointer Uniform %uint
+     %buffer = OpVariable %ptr_buffer_type Uniform
+
+       %main = OpFunction %void None %15
+         %19 = OpLabel
+               ; Access chain into atomic
+         %22 = OpAccessChain %ptr_uint %buffer %int_0 %int_0
+               ; Load of the buffer value as a u32
+         %23 = OpLoad %uint %22
+
+               ; If
+               OpSelectionMerge %25 None
+               OpBranchConditional %true %26 %25
+
+               ; True block
+         %26 = OpLabel
+               ; Access atomic, attempts to makes %22 an atomic value
+         %27 = OpAtomicCompareExchange %uint %22 %uint_1 %uint_0 %uint_0 %uint_0 %uint_0
+               OpBranch %25
+
+               ; If merge/False block
+         %25 = OpLabel
+               OpReturn
+               OpFunctionEnd