[spirv-reader][ir] Add support for OpAtomicCompareExchange.
Convert the `OpAtomicCompareExchange` instruction into an
`atomicCompareExchangeWeak` in WGSL. The result of the original
expression needs to be extracted and used as the builtin result.
Bug: 391487024
Change-Id: I42c22da639340776dded0094e0bc73c39e5b8522
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/236915
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/reader/lower/atomics.cc b/src/tint/lang/spirv/reader/lower/atomics.cc
index 8409ee1..0d76df2 100644
--- a/src/tint/lang/spirv/reader/lower/atomics.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics.cc
@@ -32,6 +32,7 @@
#include "src/tint/lang/core/ir/builder.h"
#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/core/type/builtin_structs.h"
#include "src/tint/lang/spirv/ir/builtin_call.h"
#include "src/tint/utils/containers/hashmap.h"
#include "src/tint/utils/containers/hashset.h"
@@ -91,6 +92,7 @@
AtomicOp(builtin, core::BuiltinFn::kAtomicExchange);
break;
case spirv::BuiltinFn::kAtomicCompareExchange:
+ AtomicCompareExchange(builtin);
break;
case spirv::BuiltinFn::kAtomicIAdd:
AtomicOp(builtin, core::BuiltinFn::kAtomicAdd);
@@ -150,6 +152,25 @@
TINT_ICE_ON_NO_MATCH);
}
+ void AtomicCompareExchange(spirv::ir::BuiltinCall* call) {
+ auto args = call->Args();
+
+ b.InsertBefore(call, [&] {
+ auto* var = args[0];
+ values_to_convert_.Push(var);
+
+ auto* val = args[4];
+ auto* comp = args[5];
+
+ auto* strct =
+ core::type::CreateAtomicCompareExchangeResult(ty, ir.symbols, val->Type());
+
+ auto* bi = b.Call(strct, core::BuiltinFn::kAtomicCompareExchangeWeak, var, val, comp);
+ b.AccessWithResult(call->DetachResult(), bi, 0_u);
+ });
+ call->Destroy();
+ }
+
void AtomicChangeByOne(spirv::ir::BuiltinCall* call, core::BuiltinFn fn) {
auto args = call->Args();
diff --git a/src/tint/lang/spirv/reader/lower/atomics_test.cc b/src/tint/lang/spirv/reader/lower/atomics_test.cc
index 1de66a5..b36568b 100644
--- a/src/tint/lang/spirv/reader/lower/atomics_test.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics_test.cc
@@ -1497,7 +1497,7 @@
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicCompareExchange) {
+TEST_F(SpirvReader_AtomicsTest, AtomicCompareExchange) {
auto* f = b.ComputeFunction("main");
auto* sb = ty.Struct(mod.symbols.New("S"), {
@@ -1519,15 +1519,15 @@
b.Append(f->Block(), [&] { //
auto* a0 = b.Access(ty.ptr<storage, i32, read_write>(), sg, 0_u);
b.Call<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kAtomicCompareExchange, a0, 1_u,
- 0_u, 0_u, 1_i, 1_i);
+ 0_u, 0_u, 2_i, 3_i);
auto* a1 = b.Access(ty.ptr<storage, u32, read_write>(), sg, 1_u);
b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicCompareExchange, a1, 1_u,
- 0_u, 0_u, 1_u, 1_u);
+ 0_u, 0_u, 4_u, 5_u);
b.Call<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kAtomicCompareExchange, wg_i32,
- 1_u, 0_u, 0_u, 1_i, 1_i);
+ 1_u, 0_u, 0_u, 6_i, 7_i);
b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicCompareExchange, wg_u32,
- 1_u, 0_u, 0_u, 1_u, 1_u);
+ 1_u, 0_u, 0_u, 8_u, 9_u);
b.Return(f);
});
@@ -1546,11 +1546,11 @@
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B2: {
%5:ptr<storage, i32, read_write> = access %sb, 0u
- %6:i32 = spirv.atomic_compare_exchange %5, 1u, 0u, 0u, 1i, 1i
+ %6:i32 = spirv.atomic_compare_exchange %5, 1u, 0u, 0u, 2i, 3i
%7:ptr<storage, u32, read_write> = access %sb, 1u
- %8:u32 = spirv.atomic_compare_exchange %7, 1u, 0u, 0u, 1u, 1u
- %9:i32 = spirv.atomic_compare_exchange %wg_i32, 1u, 0u, 0u, 1i, 1i
- %10:u32 = spirv.atomic_compare_exchange %wg_u32, 1u, 0u, 0u, 1u, 1u
+ %8:u32 = spirv.atomic_compare_exchange %7, 1u, 0u, 0u, 4u, 5u
+ %9:i32 = spirv.atomic_compare_exchange %wg_i32, 1u, 0u, 0u, 6i, 7i
+ %10:u32 = spirv.atomic_compare_exchange %wg_u32, 1u, 0u, 0u, 8u, 9u
ret
}
}
@@ -1560,7 +1560,47 @@
Run(Atomics);
auto* expect = R"(
-UNIMPLEMENTED
+S = struct @align(4) {
+ a:i32 @offset(0)
+ b:u32 @offset(4)
+}
+
+__atomic_compare_exchange_result_i32 = struct @align(4) {
+ old_value:i32 @offset(0)
+ exchanged:bool @offset(4)
+}
+
+__atomic_compare_exchange_result_u32 = struct @align(4) {
+ old_value:u32 @offset(0)
+ exchanged:bool @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:__atomic_compare_exchange_result_i32 = atomicCompareExchangeWeak %5, 2i, 3i
+ %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
+ %10:u32 = access %9, 0u
+ %11:__atomic_compare_exchange_result_i32 = atomicCompareExchangeWeak %wg_i32, 6i, 7i
+ %12:i32 = access %11, 0u
+ %13:__atomic_compare_exchange_result_u32 = atomicCompareExchangeWeak %wg_u32, 8u, 9u
+ %14:u32 = access %13, 0u
+ ret
+ }
+}
)";
ASSERT_EQ(expect, str());
}
