[spirv-reader][ir] Convert several spirv atomic methods.
Convert the spirv atomic methods with the same form to the equivalent
WGSL builtins.
Bug: 391487374, 391487608, 391486895, 391487060, 391486815, 391487220
Bug: 391486699, 391487749
Change-Id: I7ddd60acf5d585bf04d7472640e9ca49a0c9005a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/233236
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 9b8db28..244d8ea 100644
--- a/src/tint/lang/spirv/reader/lower/atomics.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics.cc
@@ -85,22 +85,38 @@
case spirv::BuiltinFn::kAtomicLoad:
break;
case spirv::BuiltinFn::kAtomicStore:
- AtomicStore(builtin);
+ AtomicOp(builtin, core::BuiltinFn::kAtomicStore);
break;
case spirv::BuiltinFn::kAtomicExchange:
case spirv::BuiltinFn::kAtomicCompareExchange:
break;
case spirv::BuiltinFn::kAtomicIAdd:
- AtomicIAdd(builtin);
+ AtomicOp(builtin, core::BuiltinFn::kAtomicAdd);
break;
case spirv::BuiltinFn::kAtomicISub:
+ AtomicOp(builtin, core::BuiltinFn::kAtomicSub);
+ break;
case spirv::BuiltinFn::kAtomicSMax:
+ AtomicOp(builtin, core::BuiltinFn::kAtomicMax);
+ break;
case spirv::BuiltinFn::kAtomicSMin:
+ AtomicOp(builtin, core::BuiltinFn::kAtomicMin);
+ break;
case spirv::BuiltinFn::kAtomicUMax:
+ AtomicOp(builtin, core::BuiltinFn::kAtomicMax);
+ break;
case spirv::BuiltinFn::kAtomicUMin:
+ AtomicOp(builtin, core::BuiltinFn::kAtomicMin);
+ break;
case spirv::BuiltinFn::kAtomicAnd:
+ AtomicOp(builtin, core::BuiltinFn::kAtomicAnd);
+ break;
case spirv::BuiltinFn::kAtomicOr:
+ AtomicOp(builtin, core::BuiltinFn::kAtomicOr);
+ break;
case spirv::BuiltinFn::kAtomicXor:
+ AtomicOp(builtin, core::BuiltinFn::kAtomicXor);
+ break;
case spirv::BuiltinFn::kAtomicIIncrement:
case spirv::BuiltinFn::kAtomicIDecrement:
break;
@@ -121,7 +137,7 @@
ReplaceStructTypes();
}
- void AtomicIAdd(spirv::ir::BuiltinCall* call) {
+ void AtomicOp(spirv::ir::BuiltinCall* call, core::BuiltinFn fn) {
auto args = call->Args();
b.InsertBefore(call, [&] {
@@ -129,20 +145,7 @@
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();
-
- b.InsertBefore(call, [&] {
- auto* var = args[0];
- values_to_convert_.Push(var);
-
- auto* val = args[3];
- b.CallWithResult(call->DetachResult(), core::BuiltinFn::kAtomicStore, var, val);
+ b.CallWithResult(call->DetachResult(), fn, var, val);
});
call->Destroy();
}
diff --git a/src/tint/lang/spirv/reader/lower/atomics_test.cc b/src/tint/lang/spirv/reader/lower/atomics_test.cc
index e51b861..95a82a5 100644
--- a/src/tint/lang/spirv/reader/lower/atomics_test.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics_test.cc
@@ -783,7 +783,7 @@
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicSub) {
+TEST_F(SpirvReader_AtomicsTest, AtomicSub) {
auto* f = b.ComputeFunction("main");
auto* sb = ty.Struct(mod.symbols.New("S"), {
@@ -844,12 +844,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_AtomicAnd) {
+TEST_F(SpirvReader_AtomicsTest, AtomicAnd) {
auto* f = b.ComputeFunction("main");
auto* sb = ty.Struct(mod.symbols.New("S"), {
@@ -910,12 +936,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 = atomicAnd %5, 1i
+ %7:ptr<storage, atomic<u32>, read_write> = access %sb, 1u
+ %8:u32 = atomicAnd %7, 1u
+ %9:i32 = atomicAnd %wg_i32, 1i
+ %10:u32 = atomicAnd %wg_u32, 1u
+ ret
+ }
+}
)";
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicOr) {
+TEST_F(SpirvReader_AtomicsTest, AtomicOr) {
auto* f = b.ComputeFunction("main");
auto* sb = ty.Struct(mod.symbols.New("S"), {
@@ -975,12 +1027,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 = atomicOr %5, 1i
+ %7:ptr<storage, atomic<u32>, read_write> = access %sb, 1u
+ %8:u32 = atomicOr %7, 1u
+ %9:i32 = atomicOr %wg_i32, 1i
+ %10:u32 = atomicOr %wg_u32, 1u
+ ret
+ }
+}
)";
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicXor) {
+TEST_F(SpirvReader_AtomicsTest, AtomicXor) {
auto* f = b.ComputeFunction("main");
auto* sb = ty.Struct(mod.symbols.New("S"), {
@@ -1041,12 +1119,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 = atomicXor %5, 1i
+ %7:ptr<storage, atomic<u32>, read_write> = access %sb, 1u
+ %8:u32 = atomicXor %7, 1u
+ %9:i32 = atomicXor %wg_i32, 1i
+ %10:u32 = atomicXor %wg_u32, 1u
+ ret
+ }
+}
)";
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicMax) {
+TEST_F(SpirvReader_AtomicsTest, AtomicMax) {
auto* f = b.ComputeFunction("main");
auto* sb = ty.Struct(mod.symbols.New("S"), {
@@ -1107,12 +1211,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 = atomicMax %5, 1i
+ %7:ptr<storage, atomic<u32>, read_write> = access %sb, 1u
+ %8:u32 = atomicMax %7, 1u
+ %9:i32 = atomicMax %wg_i32, 1i
+ %10:u32 = atomicMax %wg_u32, 1u
+ ret
+ }
+}
)";
ASSERT_EQ(expect, str());
}
-TEST_F(SpirvReader_AtomicsTest, DISABLED_AtomicMin) {
+TEST_F(SpirvReader_AtomicsTest, AtomicMin) {
auto* f = b.ComputeFunction("main");
auto* sb = ty.Struct(mod.symbols.New("S"), {
@@ -1173,7 +1303,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 = atomicMin %5, 1i
+ %7:ptr<storage, atomic<u32>, read_write> = access %sb, 1u
+ %8:u32 = atomicMin %7, 1u
+ %9:i32 = atomicMin %wg_i32, 1i
+ %10:u32 = atomicMin %wg_u32, 1u
+ ret
+ }
+}
)";
ASSERT_EQ(expect, str());
}
