commit f850dcb9b48c5122ee57d44f5b6c16e761414e6a
author dan sinclair <dsinclair@chromium.org> Wed Mar 26 20:41:59 2025 -0700
committer Dawn LUCI CQ <dawn-scoped@luci-project-accounts.iam.gserviceaccount.com> Wed Mar 26 20:41:59 2025 -0700
tree 9649fb776a42daf3a6ee6f5a93be36796f7f533e
parent b2da119bd31b8509f4e7a44043612bf2c9c2db7e

[spirv-reader][ir] Support spirv.atomic_i_add

Add conversion of the `spirv.atomic_i_add` instruction to an `atomicAdd`.

Bug: 391487510
Change-Id: I6b556dfa033df42ef7c28bfa9363e57343df8f9c
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/233234
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: 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 15e3379..9b8db28 100644
--- a/src/tint/lang/spirv/reader/lower/atomics.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics.cc
@@ -89,7 +89,10 @@
                     break;
                 case spirv::BuiltinFn::kAtomicExchange:
                 case spirv::BuiltinFn::kAtomicCompareExchange:
+                    break;
                 case spirv::BuiltinFn::kAtomicIAdd:
+                    AtomicIAdd(builtin);
+                    break;
                 case spirv::BuiltinFn::kAtomicISub:
                 case spirv::BuiltinFn::kAtomicSMax:
                 case spirv::BuiltinFn::kAtomicSMin:
@@ -118,6 +121,19 @@
         ReplaceStructTypes();
     }
 
+    void AtomicIAdd(spirv::ir::BuiltinCall* call) {
+        auto args = call->Args();
+
+        b.InsertBefore(call, [&] {
+            auto* var = args[0];
+            values_to_convert_.Push(var);
+
+            auto* val = args[3];
+            b.CallWithResult(call->DetachResult(), core::BuiltinFn::kAtomicAdd, var, val);
+        });
+        call->Destroy();
+    }
+
     void AtomicStore(spirv::ir::BuiltinCall* call) {
         auto args = call->Args();
 

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 12b01d3..e51b861 100644
--- a/src/tint/lang/spirv/reader/lower/atomics_test.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics_test.cc
@@ -691,7 +691,7 @@
     ASSERT_EQ(expect, str());
 }
 
-TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicAdd) {
+TEST_F(SpirvReader_AtomicsTest, AtomicAdd) {
     auto* f = b.ComputeFunction("main");
 
     auto* sb = ty.Struct(mod.symbols.New("S"), {
@@ -752,7 +752,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:i32 = atomicAdd %5, 1i
+    %7:ptr<storage, atomic<u32>, read_write> = access %sb, 1u
+    %8:u32 = atomicAdd %7, 1u
+    %9:i32 = atomicAdd %wg_i32, 1i
+    %10:u32 = atomicAdd %wg_u32, 1u
+    ret
+  }
+}
 )";
     ASSERT_EQ(expect, str());
 }