writer/wgsl: Emit access mode on pointer types
Change-Id: If694489a079698df7d767967898d6c5238fe9f54
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59821
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc
index d94299d..3dc824d 100644
--- a/src/writer/wgsl/generator_impl.cc
+++ b/src/writer/wgsl/generator_impl.cc
@@ -420,6 +420,12 @@
if (!EmitType(out, ptr->type())) {
return false;
}
+ if (ptr->access() != ast::Access::kUndefined) {
+ out << ", ";
+ if (!EmitAccess(out, ptr->access())) {
+ return false;
+ }
+ }
out << ">";
} else if (auto* atomic = ty->As<ast::Atomic>()) {
out << "atomic<";
diff --git a/src/writer/wgsl/generator_impl_type_test.cc b/src/writer/wgsl/generator_impl_type_test.cc
index 3c4d4c3..ba8e599 100644
--- a/src/writer/wgsl/generator_impl_type_test.cc
+++ b/src/writer/wgsl/generator_impl_type_test.cc
@@ -125,6 +125,18 @@
EXPECT_EQ(out.str(), "ptr<workgroup, f32>");
}
+TEST_F(WgslGeneratorImplTest, EmitType_PointerAccessMode) {
+ auto* p =
+ ty.pointer<f32>(ast::StorageClass::kStorage, ast::Access::kReadWrite);
+ Alias("make_type_reachable", p);
+
+ GeneratorImpl& gen = Build();
+
+ std::stringstream out;
+ ASSERT_TRUE(gen.EmitType(out, p)) << gen.error();
+ EXPECT_EQ(out.str(), "ptr<storage, f32, read_write>");
+}
+
TEST_F(WgslGeneratorImplTest, EmitType_Struct) {
auto* s = Structure("S", {
Member("a", ty.i32()),
diff --git a/test/bug/tint/221.wgsl.expected.wgsl b/test/bug/tint/221.wgsl.expected.wgsl
index 806a365..70d478f 100644
--- a/test/bug/tint/221.wgsl.expected.wgsl
+++ b/test/bug/tint/221.wgsl.expected.wgsl
@@ -15,7 +15,7 @@
if ((i >= b.count)) {
break;
}
- let p : ptr<storage, u32> = &(b.data[i]);
+ let p : ptr<storage, u32, read_write> = &(b.data[i]);
if (((i % 2u) == 0u)) {
continue;
}
diff --git a/test/bug/tint/492.wgsl.expected.wgsl b/test/bug/tint/492.wgsl.expected.wgsl
index 07746b2..6245430 100644
--- a/test/bug/tint/492.wgsl.expected.wgsl
+++ b/test/bug/tint/492.wgsl.expected.wgsl
@@ -7,6 +7,6 @@
[[stage(compute), workgroup_size(1)]]
fn main() {
- let p : ptr<storage, i32> = &(buf.a);
+ let p : ptr<storage, i32, read_write> = &(buf.a);
*(p) = 12;
}
diff --git a/test/ptr_ref/load/local/ptr_storage.wgsl.expected.wgsl b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.wgsl
index 65ccd81..b88f09f 100644
--- a/test/ptr_ref/load/local/ptr_storage.wgsl.expected.wgsl
+++ b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.wgsl
@@ -7,6 +7,6 @@
[[stage(compute), workgroup_size(1)]]
fn main() {
- let p : ptr<storage, i32> = &(v.a);
+ let p : ptr<storage, i32, read_write> = &(v.a);
let use : i32 = (*(p) + 1);
}
