From 3d73768aec3e84fe58a14a526e7fb4a0b1f80a75 Mon Sep 17 00:00:00 2001
From: James Price <jrprice@google.com>
Date: Mon, 26 Jul 2021 22:11:58 +0000
Subject: [PATCH] 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                    |  6 ++++++
 src/writer/wgsl/generator_impl_type_test.cc          | 12 ++++++++++++
 test/bug/tint/221.wgsl.expected.wgsl                 |  2 +-
 test/bug/tint/492.wgsl.expected.wgsl                 |  2 +-
 .../load/local/ptr_storage.wgsl.expected.wgsl        |  2 +-
 5 files changed, 21 insertions(+), 3 deletions(-)

diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc
index d94299d8d2..3dc824dc91 100644
--- a/src/writer/wgsl/generator_impl.cc
+++ b/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<";
diff --git a/src/writer/wgsl/generator_impl_type_test.cc b/src/writer/wgsl/generator_impl_type_test.cc
index 3c4d4c3e43..ba8e599e85 100644
--- a/src/writer/wgsl/generator_impl_type_test.cc
+++ b/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()),
diff --git a/test/bug/tint/221.wgsl.expected.wgsl b/test/bug/tint/221.wgsl.expected.wgsl
index 806a365820..70d478f47a 100644
--- a/test/bug/tint/221.wgsl.expected.wgsl
+++ b/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;
     }
diff --git a/test/bug/tint/492.wgsl.expected.wgsl b/test/bug/tint/492.wgsl.expected.wgsl
index 07746b24eb..6245430c17 100644
--- a/test/bug/tint/492.wgsl.expected.wgsl
+++ b/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;
 }
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 65ccd81e65..b88f09f6ea 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 @@ 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);
 }