[spirv-reader][ir] Support atomic increment and decrement.
Add support to translate the SPIR-V atomic_i_increment and
atomic_i_decrement instructions to the WGSL atomic{add,Sub} operations
which we use to polyfill.
Bug: 391486942, 391487231
Change-Id: I43bf1c478619d7f9ca2cc1527061c2a80f9430b2
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/236194
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Auto-Submit: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/lower/atomics.cc b/src/tint/lang/spirv/reader/lower/atomics.cc
index 7c422ea..2e063cb 100644
--- a/src/tint/lang/spirv/reader/lower/atomics.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics.cc
@@ -118,7 +118,10 @@
AtomicOp(builtin, core::BuiltinFn::kAtomicXor);
break;
case spirv::BuiltinFn::kAtomicIIncrement:
+ AtomicChangeByOne(builtin, core::BuiltinFn::kAtomicAdd);
+ break;
case spirv::BuiltinFn::kAtomicIDecrement:
+ AtomicChangeByOne(builtin, core::BuiltinFn::kAtomicSub);
break;
default:
TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
@@ -137,6 +140,27 @@
ReplaceStructTypes();
}
+ core::ir::Value* One(const core::type::Type* const_ty) {
+ return tint::Switch(
+ const_ty, //
+ [&](const core::type::I32*) { return b.Constant(1_i); },
+ [&](const core::type::U32*) { return b.Constant(1_u); }, //
+ TINT_ICE_ON_NO_MATCH);
+ }
+
+ void AtomicChangeByOne(spirv::ir::BuiltinCall* call, core::BuiltinFn fn) {
+ auto args = call->Args();
+
+ b.InsertBefore(call, [&] {
+ auto* var = args[0];
+ values_to_convert_.Push(var);
+
+ auto* one = One(call->Result()->Type());
+ b.CallWithResult(call->DetachResult(), fn, var, one);
+ });
+ call->Destroy();
+ }
+
void AtomicOp(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 0bf3fa7..c7f21a6 100644
--- a/src/tint/lang/spirv/reader/lower/atomics_test.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics_test.cc
@@ -1671,7 +1671,7 @@
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicDecrement) {
+TEST_F(SpirvReader_AtomicsTest, AtomicDecrement) {
auto* f = b.ComputeFunction("main");
auto* sb = ty.Struct(mod.symbols.New("S"), {
@@ -1695,10 +1695,10 @@
b.Call<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kAtomicIDecrement, a0, 1_u, 0_u);
auto* a1 = b.Access(ty.ptr<storage, u32, read_write>(), sg, 1_u);
- b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicIDecrement, a1, 1_u, 0_u);
- b.Call<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kAtomicIDecrement, wg_i32, 1_u,
+ b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicIDecrement, a1, 4_u, 0_u);
+ b.Call<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kAtomicIDecrement, wg_i32, 3_u,
0_u);
- b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicIDecrement, wg_u32, 1_u,
+ b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicIDecrement, wg_u32, 2_u,
0_u);
b.Return(f);
});
@@ -1720,9 +1720,9 @@
%5:ptr<storage, i32, read_write> = access %sb, 0u
%6:i32 = spirv.atomic_i_decrement %5, 1u, 0u
%7:ptr<storage, u32, read_write> = access %sb, 1u
- %8:u32 = spirv.atomic_i_decrement %7, 1u, 0u
- %9:i32 = spirv.atomic_i_decrement %wg_i32, 1u, 0u
- %10:u32 = spirv.atomic_i_decrement %wg_u32, 1u, 0u
+ %8:u32 = spirv.atomic_i_decrement %7, 4u, 0u
+ %9:i32 = spirv.atomic_i_decrement %wg_i32, 3u, 0u
+ %10:u32 = spirv.atomic_i_decrement %wg_u32, 2u, 0u
ret
}
}
@@ -1732,12 +1732,38 @@
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 = atomicSub %5, 1i
+ %7:ptr<storage, atomic<u32>, read_write> = access %sb, 1u
+ %8:u32 = atomicSub %7, 1u
+ %9:i32 = atomicSub %wg_i32, 1i
+ %10:u32 = atomicSub %wg_u32, 1u
+ ret
+ }
+}
)";
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicIncrement) {
+TEST_F(SpirvReader_AtomicsTest, AtomicIncrement) {
auto* f = b.ComputeFunction("main");
auto* sb = ty.Struct(mod.symbols.New("S"), {
@@ -1761,10 +1787,10 @@
b.Call<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kAtomicIIncrement, a0, 1_u, 0_u);
auto* a1 = b.Access(ty.ptr<storage, u32, read_write>(), sg, 1_u);
- b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicIIncrement, a1, 1_u, 0_u);
- b.Call<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kAtomicIIncrement, wg_i32, 1_u,
+ b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicIIncrement, a1, 4_u, 0_u);
+ b.Call<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kAtomicIIncrement, wg_i32, 3_u,
0_u);
- b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicIIncrement, wg_u32, 1_u,
+ b.Call<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kAtomicIIncrement, wg_u32, 2_u,
0_u);
b.Return(f);
});
@@ -1786,9 +1812,9 @@
%5:ptr<storage, i32, read_write> = access %sb, 0u
%6:i32 = spirv.atomic_i_increment %5, 1u, 0u
%7:ptr<storage, u32, read_write> = access %sb, 1u
- %8:u32 = spirv.atomic_i_increment %7, 1u, 0u
- %9:i32 = spirv.atomic_i_increment %wg_i32, 1u, 0u
- %10:u32 = spirv.atomic_i_increment %wg_u32, 1u, 0u
+ %8:u32 = spirv.atomic_i_increment %7, 4u, 0u
+ %9:i32 = spirv.atomic_i_increment %wg_i32, 3u, 0u
+ %10:u32 = spirv.atomic_i_increment %wg_u32, 2u, 0u
ret
}
}
@@ -1798,7 +1824,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());
}
