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>

src/writer/wgsl/generator_impl.cc

@@ -420,6 +420,12 @@ bool GeneratorImpl::EmitType(std::ostream& out, const ast::Type* ty) {
     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<";

src/writer/wgsl/generator_impl_type_test.cc

@@ -125,6 +125,18 @@ TEST_F(WgslGeneratorImplTest, EmitType_Pointer) {
   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()),

test/bug/tint/221.wgsl.expected.wgsl

@@ -15,7 +15,7 @@ fn main() {
     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;
     }

test/bug/tint/492.wgsl.expected.wgsl

@@ -7,6 +7,6 @@ struct S {
 
 [[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;
 }

test/ptr_ref/load/local/ptr_storage.wgsl.expected.wgsl

@@ -7,6 +7,6 @@ struct S {
 
 [[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);
 }