[spirv-reader] Allow arrayLength on atomic arrays
This is the only other core builtin that we should see for usages of
pointers that have been converted in the Atomics transform.
Fixed: 435502510
Change-Id: Ia8af21d74b33246056a19da97af83ec029ee1a40
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/255559
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: 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 7a64995..0820b0d 100644
--- a/src/tint/lang/spirv/reader/lower/atomics.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics.cc
@@ -399,6 +399,12 @@
},
[&](core::ir::UserCall* uc) { user_calls_to_convert_.Add(uc); },
[&](core::ir::CoreBuiltinCall* bc) {
+ // The only non-atomic builtin that can touch something that contains an atomic
+ // is arrayLength.
+ if (bc->Func() == core::BuiltinFn::kArrayLength) {
+ return;
+ }
+
// This was converted when we switched from a SPIR-V intrinsic to core
TINT_ASSERT(core::IsAtomic(bc->Func()));
TINT_ASSERT(bc->Args()[0]->Type()->UnwrapPtr()->Is<core::type::Atomic>());
diff --git a/src/tint/lang/spirv/reader/lower/atomics_test.cc b/src/tint/lang/spirv/reader/lower/atomics_test.cc
index 47072da..50d1c76 100644
--- a/src/tint/lang/spirv/reader/lower/atomics_test.cc
+++ b/src/tint/lang/spirv/reader/lower/atomics_test.cc
@@ -3203,5 +3203,57 @@
ASSERT_EQ(expect, str());
}
+TEST_F(SpirvReader_AtomicsTest, ArrayLength) {
+ auto* f = b.ComputeFunction("main");
+
+ core::ir::Var* buffer = nullptr;
+ b.Append(mod.root_block,
+ [&] { //
+ buffer = b.Var("buffer", ty.ptr<storage, array<u32>, read_write>());
+ buffer->SetBindingPoint(0u, 0u);
+ });
+
+ b.Append(f->Block(), [&] { //
+ auto* a = b.Access(ty.ptr<storage, u32, read_write>(), buffer, 1_i);
+ b.Call<spirv::ir::BuiltinCall>(ty.void_(), spirv::BuiltinFn::kAtomicStore, a, 1_u, 0_u,
+ 1_u);
+ b.Call<core::ir::CoreBuiltinCall>(ty.u32(), core::BuiltinFn::kArrayLength, buffer);
+ b.Return(f);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %buffer:ptr<storage, array<u32>, read_write> = var undef @binding_point(0, 0)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %3:ptr<storage, u32, read_write> = access %buffer, 1i
+ %4:void = spirv.atomic_store %3, 1u, 0u, 1u
+ %5:u32 = arrayLength %buffer
+ ret
+ }
+}
+)";
+ ASSERT_EQ(src, str());
+ Run(Atomics);
+
+ auto* expect = R"(
+$B1: { # root
+ %buffer:ptr<storage, array<atomic<u32>>, read_write> = var undef @binding_point(0, 0)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %3:ptr<storage, atomic<u32>, read_write> = access %buffer, 1i
+ %4:void = atomicStore %3, 1u
+ %5:u32 = arrayLength %buffer
+ ret
+ }
+}
+)";
+ ASSERT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::spirv::reader::lower
