[spirv-reader][ir] Convert OpAtomicExchange

Convert the `OpAtomicExchange` to an `atomicExchange` instruction in
WGSL.

Bug: 391486936
Change-Id: Ie0e08cff353e6e4308d3c765341a7fd3314a8dc5
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/236895
Commit-Queue: dan sinclair <[email protected]>
Reviewed-by: James Price <[email protected]>

Author: dj2

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

@@ -88,6 +88,8 @@ struct State {
                     AtomicOp(builtin, core::BuiltinFn::kAtomicStore);
                     break;
                 case spirv::BuiltinFn::kAtomicExchange:
+                    AtomicOp(builtin, core::BuiltinFn::kAtomicExchange);
+                    break;
                 case spirv::BuiltinFn::kAtomicCompareExchange:
                     break;
                 case spirv::BuiltinFn::kAtomicIAdd:

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

@@ -1403,7 +1403,7 @@ S_atomic = struct @align(4) {
     ASSERT_EQ(expect, str());
 }
 
-TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicExchange) {
+TEST_F(SpirvReader_AtomicsTest, AtomicExchange) {
     auto* f = b.ComputeFunction("main");
 
     auto* sb = ty.Struct(mod.symbols.New("S"), {
@@ -1429,11 +1429,11 @@ TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicExchange) {
 
         auto* a1 = b.Access(ty.ptr<storage, u32, read_write>(), sg, 1_u);
         b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicExchange, a1, 1_u, 0_u,
-                                       1_u);
+                                       2_u);
         b.Call<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kAtomicExchange, wg_i32, 1_u,
-                                       0_u, 1_i);
+                                       0_u, 3_i);
         b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicExchange, wg_u32, 1_u,
-                                       0_u, 1_u);
+                                       0_u, 4_u);
         b.Return(f);
     });
 
@@ -1454,9 +1454,9 @@ S = struct @align(4) {
     %5:ptr<storage, i32, read_write> = access %sb, 0u
     %6:i32 = spirv.atomic_exchange %5, 1u, 0u, 1i
     %7:ptr<storage, u32, read_write> = access %sb, 1u
-    %8:u32 = spirv.atomic_exchange %7, 1u, 0u, 1u
-    %9:i32 = spirv.atomic_exchange %wg_i32, 1u, 0u, 1i
-    %10:u32 = spirv.atomic_exchange %wg_u32, 1u, 0u, 1u
+    %8:u32 = spirv.atomic_exchange %7, 1u, 0u, 2u
+    %9:i32 = spirv.atomic_exchange %wg_i32, 1u, 0u, 3i
+    %10:u32 = spirv.atomic_exchange %wg_u32, 1u, 0u, 4u
     ret
   }
 }
@@ -1466,7 +1466,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:i32 = atomicExchange %5, 1i
+    %7:ptr<storage, atomic<u32>, read_write> = access %sb, 1u
+    %8:u32 = atomicExchange %7, 2u
+    %9:i32 = atomicExchange %wg_i32, 3i
+    %10:u32 = atomicExchange %wg_u32, 4u
+    ret
+  }
+}
 )";
     ASSERT_EQ(expect, str());
 }