diff --git a/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc b/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc
index efb14a2..8e486f5 100644
--- a/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc
@@ -1534,14 +1534,14 @@
             // here.)
             ast_initializer = MakeConstantExpression(var.GetSingleWordInOperand(1)).expr;
         }
-        auto ast_access = VarAccess(ast_store_type, ast_address_space);
-        auto* ast_var = MakeVar(var.result_id(), ast_address_space, ast_access, ast_store_type,
-                                ast_initializer, Attributes{});
+        auto* ast_var = MakeVar(var.result_id(), ast_address_space, ast_store_type, ast_initializer,
+                                Attributes{});
         // TODO(dneto): initializers (a.k.a. initializer expression)
         if (ast_var) {
             builder_.AST().AddGlobalVariable(ast_var);
             module_variable_.GetOrAdd(var.result_id(), [&] {
-                return ModuleVariable{ast_var, ast_address_space, ast_access};
+                auto access = VarAccess(var.result_id(), ast_store_type, ast_address_space);
+                return ModuleVariable{ast_var, ast_address_space, access};
             });
         }
     }
@@ -1571,8 +1571,7 @@
         }
         auto storage_type = ConvertType(builtin_position_.position_member_type_id);
         auto ast_address_space = enum_converter_.ToAddressSpace(builtin_position_.storage_class);
-        auto ast_access = VarAccess(storage_type, ast_address_space);
-        auto* ast_var = MakeVar(builtin_position_.per_vertex_var_id, ast_address_space, ast_access,
+        auto* ast_var = MakeVar(builtin_position_.per_vertex_var_id, ast_address_space,
                                 storage_type, ast_initializer, {});
 
         builder_.AST().AddGlobalVariable(ast_var);
@@ -1606,14 +1605,16 @@
     return size->AsIntConstant();
 }
 
-core::Access ASTParser::VarAccess(const Type* storage_type, core::AddressSpace address_space) {
+core::Access ASTParser::VarAccess(uint32_t var_id,
+                                  const Type* storage_type,
+                                  core::AddressSpace address_space) {
     if (address_space != core::AddressSpace::kStorage) {
         return core::Access::kUndefined;
     }
 
-    bool read_only = false;
+    bool read_only = read_only_vars_.count(var_id) > 0;
     if (auto* tn = storage_type->As<Named>()) {
-        read_only = read_only_struct_types_.count(tn->name) > 0;
+        read_only = read_only || read_only_struct_types_.count(tn->name) > 0;
     }
 
     // Apply the access(read) or access(read_write) modifier.
@@ -1622,7 +1623,6 @@
 
 const ast::Var* ASTParser::MakeVar(uint32_t id,
                                    core::AddressSpace address_space,
-                                   core::Access access,
                                    const Type* storage_type,
                                    const ast::Expression* initializer,
                                    Attributes attrs) {
@@ -1642,6 +1642,8 @@
         return nullptr;
     }
 
+    const auto access = VarAccess(id, storage_type, address_space);
+
     // Use type inference if there is an initializer.
     auto sym = builder_.Symbols().Register(namer_.Name(id));
     return builder_.Var(Source{}, sym, initializer ? ast::Type{} : storage_type->Build(builder_),
@@ -1755,6 +1757,9 @@
             }
             attrs.Add(builder_.Binding(Source{}, AInt(deco[1])));
         }
+        if (deco[0] == uint32_t(spv::Decoration::NonWritable)) {
+            read_only_vars_.insert(id);
+        }
     }
 
     if (transfer_pipeline_io) {
diff --git a/src/tint/lang/spirv/reader/ast_parser/ast_parser.h b/src/tint/lang/spirv/reader/ast_parser/ast_parser.h
index 106221a..4174965 100644
--- a/src/tint/lang/spirv/reader/ast_parser/ast_parser.h
+++ b/src/tint/lang/spirv/reader/ast_parser/ast_parser.h
@@ -267,6 +267,7 @@
     /// then the `type` parameter is updated.  Returns false on failure (with
     /// a diagnostic), or when the variable should not be emitted, e.g. for a
     /// PointSize builtin.
+    /// This method is idempotent.
     /// @param id the ID of the SPIR-V variable
     /// @param store_type the WGSL store type for the variable, which should be prepopulated
     /// @param attributes the attribute list to populate
@@ -434,17 +435,20 @@
     /// @returns a list of SPIR-V decorations.
     DecorationList GetMemberPipelineDecorations(const Struct& struct_type, int member_index);
 
+    /// @param var_id the SPIR-V ID of the OpVariable
     /// @param storage_type the 'var' storage type
     /// @param address_space the 'var' address space
-    /// @returns the access mode for a 'var' declaration with the given storage type and address
-    /// space.
-    core::Access VarAccess(const Type* storage_type, core::AddressSpace address_space);
+    /// @returns the access mode for a 'var' declaration with the given variable id, storage type
+    /// and address space. Must only be called after decorations for the variable have been
+    /// converted.
+    core::Access VarAccess(uint32_t var_id,
+                           const Type* storage_type,
+                           core::AddressSpace address_space);
 
     /// Creates an AST 'var' node for a SPIR-V ID, including any attached decorations, unless it's
     /// an ignorable builtin variable.
     /// @param id the SPIR-V result ID
     /// @param address_space the address space, which cannot be core::AddressSpace::kUndefined
-    /// @param access the access
     /// @param storage_type the storage type of the variable
     /// @param initializer the variable initializer
     /// @param attributes the variable attributes
@@ -452,7 +456,6 @@
     /// in the error case
     const ast::Var* MakeVar(uint32_t id,
                             core::AddressSpace address_space,
-                            core::Access access,
                             const Type* storage_type,
                             const ast::Expression* initializer,
                             Attributes attributes);
@@ -906,6 +909,9 @@
     // The ast::Struct type names with only read-only members.
     std::unordered_set<Symbol> read_only_struct_types_;
 
+    // The IDs of variables marked as NonWritable.
+    std::unordered_set<uint32_t> read_only_vars_;
+
     // Maps from OpConstantComposite IDs to identifiers of module-scope const declarations.
     std::unordered_map<uint32_t, Symbol> declared_constant_composites_;
 
diff --git a/src/tint/lang/spirv/reader/ast_parser/function.cc b/src/tint/lang/spirv/reader/ast_parser/function.cc
index b43dc3d..9eaea1e 100644
--- a/src/tint/lang/spirv/reader/ast_parser/function.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/function.cc
@@ -2522,8 +2522,7 @@
             }
         }
         auto* var = parser_impl_.MakeVar(inst.result_id(), core::AddressSpace::kUndefined,
-                                         core::Access::kUndefined, var_store_type, initializer,
-                                         Attributes{});
+                                         var_store_type, initializer, Attributes{});
         auto* var_decl_stmt = create<ast::VariableDeclStatement>(Source{}, var);
         AddStatement(var_decl_stmt);
         auto* var_type = ty_.Reference(core::AddressSpace::kUndefined, var_store_type);
@@ -3369,9 +3368,8 @@
         // no need to remap pointer properties.
         auto* store_type = parser_impl_.ConvertType(def_inst->type_id());
         AddStatement(create<ast::VariableDeclStatement>(
-            Source{},
-            parser_impl_.MakeVar(id, core::AddressSpace::kUndefined, core::Access::kUndefined,
-                                 store_type, nullptr, Attributes{})));
+            Source{}, parser_impl_.MakeVar(id, core::AddressSpace::kUndefined, store_type, nullptr,
+                                           Attributes{})));
         auto* type = ty_.Reference(core::AddressSpace::kUndefined, store_type);
         identifier_types_.emplace(id, type);
     }
diff --git a/src/tint/lang/spirv/reader/ast_parser/module_var_test.cc b/src/tint/lang/spirv/reader/ast_parser/module_var_test.cc
index 3d67797..e3f24e1 100644
--- a/src/tint/lang/spirv/reader/ast_parser/module_var_test.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/module_var_test.cc
@@ -1394,6 +1394,37 @@
         << p->error();
 }
 
+TEST_F(SpvModuleScopeVarParserTest, StorageBuffer_NonWritable_Var) {
+    // Variable should have access(read)
+    auto p = parser(test::Assemble(Preamble() + FragMain() + R"(
+     OpDecorate %s Block
+     OpDecorate %1 DescriptorSet 0
+     OpDecorate %1 Binding 0
+     OpDecorate %1 NonWritable
+     OpMemberDecorate %s 0 Offset 0
+     OpMemberDecorate %s 1 Offset 4
+     %void = OpTypeVoid
+     %voidfn = OpTypeFunction %void
+     %float = OpTypeFloat 32
+
+     %s = OpTypeStruct %float %float
+     %ptr_sb_s = OpTypePointer StorageBuffer %s
+     %1 = OpVariable %ptr_sb_s StorageBuffer
+  )" + MainBody()));
+    ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
+    EXPECT_TRUE(p->error().empty());
+    const auto module_str = test::ToString(p->program());
+    EXPECT_THAT(module_str, HasSubstr(R"(struct S {
+  /* @offset(0) */
+  field0 : f32,
+  /* @offset(4) */
+  field1 : f32,
+}
+
+@group(0) @binding(0) var<storage, read> x_1 : S;
+)")) << module_str;
+}
+
 TEST_F(SpvModuleScopeVarParserTest, StorageBuffer_NonWritable_AllMembers) {
     // Variable should have access(read)
     auto p = parser(test::Assemble(Preamble() + FragMain() + R"(
diff --git a/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl b/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl
index 78faad9..03495e0 100644
--- a/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl
+++ b/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl
@@ -3,7 +3,7 @@
   int b;
 };
 
-RWByteAddressBuffer sb : register(u0);
+ByteAddressBuffer sb : register(t0);
 
 S sb_load(uint offset) {
   S tint_symbol = {asfloat(sb.Load4((offset + 0u))), asint(sb.Load((offset + 16u)))};
diff --git a/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl b/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl
index 78faad9..03495e0 100644
--- a/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl
+++ b/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl
@@ -3,7 +3,7 @@
   int b;
 };
 
-RWByteAddressBuffer sb : register(u0);
+ByteAddressBuffer sb : register(t0);
 
 S sb_load(uint offset) {
   S tint_symbol = {asfloat(sb.Load4((offset + 0u))), asint(sb.Load((offset + 16u)))};
diff --git a/test/tint/bug/tint/1776.spvasm.expected.msl b/test/tint/bug/tint/1776.spvasm.expected.msl
index 0933ac0..0548772 100644
--- a/test/tint/bug/tint/1776.spvasm.expected.msl
+++ b/test/tint/bug/tint/1776.spvasm.expected.msl
@@ -24,12 +24,12 @@
   /* 0x0000 */ tint_array<S, 1> inner;
 };
 
-void main_1(device sb_block* const tint_symbol_1) {
+void main_1(const device sb_block* const tint_symbol_1) {
   S const x_18 = (*(tint_symbol_1)).inner[1];
   return;
 }
 
-kernel void tint_symbol(device sb_block* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(const device sb_block* tint_symbol_2 [[buffer(0)]]) {
   main_1(tint_symbol_2);
   return;
 }
diff --git a/test/tint/bug/tint/1776.spvasm.expected.spvasm b/test/tint/bug/tint/1776.spvasm.expected.spvasm
index 6e076d1..9e220a4 100644
--- a/test/tint/bug/tint/1776.spvasm.expected.spvasm
+++ b/test/tint/bug/tint/1776.spvasm.expected.spvasm
@@ -20,6 +20,7 @@
                OpMemberDecorate %S 0 Offset 0
                OpMemberDecorate %S 1 Offset 16
                OpDecorate %_runtimearr_S ArrayStride 32
+               OpDecorate %sb NonWritable
                OpDecorate %sb DescriptorSet 0
                OpDecorate %sb Binding 0
       %float = OpTypeFloat 32
diff --git a/test/tint/bug/tint/1776.spvasm.expected.wgsl b/test/tint/bug/tint/1776.spvasm.expected.wgsl
index f9988fa..fa35f5e 100644
--- a/test/tint/bug/tint/1776.spvasm.expected.wgsl
+++ b/test/tint/bug/tint/1776.spvasm.expected.wgsl
@@ -12,7 +12,7 @@
   inner : RTArr,
 }
 
-@group(0) @binding(0) var<storage, read_write> sb : sb_block;
+@group(0) @binding(0) var<storage, read> sb : sb_block;
 
 fn main_1() {
   let x_18 = sb.inner[1i];