commit 73275c1cfcad51a4369e4db690b41c78452d5cdc
author James Price <jrprice@google.com> Fri Aug 01 17:37:15 2025 -0700
committer Dawn LUCI CQ <dawn-scoped@luci-project-accounts.iam.gserviceaccount.com> Fri Aug 01 17:37:15 2025 -0700
tree 7fee7b0c409d31753d23dc421843915418ba575f
parent 7e39f3b7f5f4639060c1206a6ddc088f17e984b7

[spirv-reader] Fix order of arguments for atomicCompareExchangeWeak

In WGSL the comparator comes first, which is the opposite to SPIR-V.

Fixed: 435495236
Change-Id: Ia1d6c8efe0aef32af26c5b869a1d143652a92ea9
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/255558
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>

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 1bcceba..7a64995 100644
--- a/src/tint/lang/spirv/reader/lower/atomics.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics.cc
@@ -221,7 +221,7 @@
             auto* strct =
                 core::type::CreateAtomicCompareExchangeResult(ty, ir.symbols, val->Type());
 
-            auto* bi = b.Call(strct, core::BuiltinFn::kAtomicCompareExchangeWeak, var, val, comp);
+            auto* bi = b.Call(strct, core::BuiltinFn::kAtomicCompareExchangeWeak, var, comp, val);
             b.AccessWithResult(call->DetachResult(), bi, 0_u);
         });
         call->Destroy();

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 a5ef867..47072da 100644
--- a/src/tint/lang/spirv/reader/lower/atomics_test.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics_test.cc
@@ -1645,14 +1645,14 @@
 %main = @compute @workgroup_size(1u, 1u, 1u) func():void {
   $B2: {
     %5:ptr<storage, atomic<i32>, read_write> = access %sb, 0u
-    %6:__atomic_compare_exchange_result_i32 = atomicCompareExchangeWeak %5, 2i, 3i
+    %6:__atomic_compare_exchange_result_i32 = atomicCompareExchangeWeak %5, 3i, 2i
     %7:i32 = access %6, 0u
     %8:ptr<storage, atomic<u32>, read_write> = access %sb, 1u
-    %9:__atomic_compare_exchange_result_u32 = atomicCompareExchangeWeak %8, 4u, 5u
+    %9:__atomic_compare_exchange_result_u32 = atomicCompareExchangeWeak %8, 5u, 4u
     %10:u32 = access %9, 0u
-    %11:__atomic_compare_exchange_result_i32 = atomicCompareExchangeWeak %wg_i32, 6i, 7i
+    %11:__atomic_compare_exchange_result_i32 = atomicCompareExchangeWeak %wg_i32, 7i, 6i
     %12:i32 = access %11, 0u
-    %13:__atomic_compare_exchange_result_u32 = atomicCompareExchangeWeak %wg_u32, 8u, 9u
+    %13:__atomic_compare_exchange_result_u32 = atomicCompareExchangeWeak %wg_u32, 9u, 8u
     %14:u32 = access %13, 0u
     ret
   }