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);
 }