diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index dd679c7..fbed83a 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1646,7 +1646,6 @@
       "reader/wgsl/classify_template_args_test.cc",
       "reader/wgsl/lexer_test.cc",
       "reader/wgsl/parser_impl_additive_expression_test.cc",
-      "reader/wgsl/parser_impl_address_space_test.cc",
       "reader/wgsl/parser_impl_argument_expression_list_test.cc",
       "reader/wgsl/parser_impl_assignment_stmt_test.cc",
       "reader/wgsl/parser_impl_bitwise_expression_test.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 67ca5e9..4b87d7b 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -1092,7 +1092,6 @@
       reader/wgsl/parser_impl_singular_expression_test.cc
       reader/wgsl/parser_impl_statement_test.cc
       reader/wgsl/parser_impl_statements_test.cc
-      reader/wgsl/parser_impl_address_space_test.cc
       reader/wgsl/parser_impl_struct_body_decl_test.cc
       reader/wgsl/parser_impl_struct_decl_test.cc
       reader/wgsl/parser_impl_struct_attribute_decl_test.cc
diff --git a/src/tint/ast/var.cc b/src/tint/ast/var.cc
index b748381..bccdbcc 100644
--- a/src/tint/ast/var.cc
+++ b/src/tint/ast/var.cc
@@ -25,8 +25,8 @@
          const Source& src,
          const Identifier* n,
          Type ty,
-         type::AddressSpace address_space,
-         type::Access access,
+         const Expression* address_space,
+         const Expression* access,
          const Expression* init,
          utils::VectorRef<const Attribute*> attrs)
     : Base(pid, nid, src, n, ty, init, std::move(attrs)),
@@ -45,10 +45,11 @@
     auto src = ctx->Clone(source);
     auto* n = ctx->Clone(name);
     auto ty = ctx->Clone(type);
+    auto* address_space = ctx->Clone(declared_address_space);
+    auto* access = ctx->Clone(declared_access);
     auto* init = ctx->Clone(initializer);
     auto attrs = ctx->Clone(attributes);
-    return ctx->dst->create<Var>(src, n, ty, declared_address_space, declared_access, init,
-                                 std::move(attrs));
+    return ctx->dst->create<Var>(src, n, ty, address_space, access, init, std::move(attrs));
 }
 
 }  // namespace tint::ast
diff --git a/src/tint/ast/var.h b/src/tint/ast/var.h
index e3537d1..006d8de 100644
--- a/src/tint/ast/var.h
+++ b/src/tint/ast/var.h
@@ -56,8 +56,8 @@
         const Source& source,
         const Identifier* name,
         Type type,
-        type::AddressSpace declared_address_space,
-        type::Access declared_access,
+        const Expression* declared_address_space,
+        const Expression* declared_access,
         const Expression* initializer,
         utils::VectorRef<const Attribute*> attributes);
 
@@ -77,10 +77,10 @@
     const Var* Clone(CloneContext* ctx) const override;
 
     /// The declared address space
-    const type::AddressSpace declared_address_space;
+    const Expression* const declared_address_space = nullptr;
 
     /// The declared access control
-    const type::Access declared_access;
+    const Expression* const declared_access = nullptr;
 };
 
 /// A list of `var` declarations
diff --git a/src/tint/ast/variable_test.cc b/src/tint/ast/variable_test.cc
index 88bcdc4..dcbf1e5 100644
--- a/src/tint/ast/variable_test.cc
+++ b/src/tint/ast/variable_test.cc
@@ -28,7 +28,8 @@
     auto* v = Var("my_var", ty.i32(), type::AddressSpace::kFunction);
 
     CheckIdentifier(Symbols(), v->name, "my_var");
-    EXPECT_EQ(v->declared_address_space, type::AddressSpace::kFunction);
+    CheckIdentifier(Symbols(), v->declared_address_space, "function");
+    EXPECT_EQ(v->declared_access, nullptr);
     CheckIdentifier(Symbols(), v->type, "i32");
     EXPECT_EQ(v->source.range.begin.line, 0u);
     EXPECT_EQ(v->source.range.begin.column, 0u);
@@ -41,7 +42,7 @@
                   ty.f32(), type::AddressSpace::kPrivate, utils::Empty);
 
     CheckIdentifier(Symbols(), v->name, "i");
-    EXPECT_EQ(v->declared_address_space, type::AddressSpace::kPrivate);
+    CheckIdentifier(Symbols(), v->declared_address_space, "private");
     CheckIdentifier(Symbols(), v->type, "f32");
     EXPECT_EQ(v->source.range.begin.line, 27u);
     EXPECT_EQ(v->source.range.begin.column, 4u);
@@ -54,7 +55,7 @@
                   ty.i32(), type::AddressSpace::kWorkgroup, utils::Empty);
 
     CheckIdentifier(Symbols(), v->name, "a_var");
-    EXPECT_EQ(v->declared_address_space, type::AddressSpace::kWorkgroup);
+    CheckIdentifier(Symbols(), v->declared_address_space, "workgroup");
     CheckIdentifier(Symbols(), v->type, "i32");
     EXPECT_EQ(v->source.range.begin.line, 27u);
     EXPECT_EQ(v->source.range.begin.column, 4u);
diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index ee28010..f55b12c 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -80,9 +80,9 @@
   uniform
   storage
   push_constant
+  __in
+  __out
   @internal handle
-  @internal in
-  @internal out
 }
 
 // https://gpuweb.github.io/gpuweb/wgsl/#memory-access-mode
diff --git a/src/tint/program_builder.h b/src/tint/program_builder.h
index 1c04faf..f89a30b 100644
--- a/src/tint/program_builder.h
+++ b/src/tint/program_builder.h
@@ -219,24 +219,34 @@
     /// constructing an ast::Var.
     struct VarOptions {
         template <typename... ARGS>
-        explicit VarOptions(ARGS&&... args) {
-            (Set(std::forward<ARGS>(args)), ...);
+        explicit VarOptions(ProgramBuilder& b, ARGS&&... args) {
+            (Set(b, std::forward<ARGS>(args)), ...);
         }
         ~VarOptions();
 
         ast::Type type;
-        type::AddressSpace address_space = type::AddressSpace::kUndefined;
-        type::Access access = type::Access::kUndefined;
+        const ast::Expression* address_space = nullptr;
+        const ast::Expression* access = nullptr;
         const ast::Expression* initializer = nullptr;
         utils::Vector<const ast::Attribute*, 4> attributes;
 
       private:
-        void Set(ast::Type t) { type = t; }
-        void Set(type::AddressSpace addr_space) { address_space = addr_space; }
-        void Set(type::Access ac) { access = ac; }
-        void Set(const ast::Expression* c) { initializer = c; }
-        void Set(utils::VectorRef<const ast::Attribute*> l) { attributes = std::move(l); }
-        void Set(const ast::Attribute* a) { attributes.Push(a); }
+        void Set(ProgramBuilder&, ast::Type t) { type = t; }
+        void Set(ProgramBuilder& b, type::AddressSpace addr_space) {
+            if (addr_space != type::AddressSpace::kUndefined) {
+                address_space = b.Expr(addr_space);
+            }
+        }
+        void Set(ProgramBuilder& b, type::Access ac) {
+            if (ac != type::Access::kUndefined) {
+                access = b.Expr(ac);
+            }
+        }
+        void Set(ProgramBuilder&, const ast::Expression* c) { initializer = c; }
+        void Set(ProgramBuilder&, utils::VectorRef<const ast::Attribute*> l) {
+            attributes = std::move(l);
+        }
+        void Set(ProgramBuilder&, const ast::Attribute* a) { attributes.Push(a); }
     };
 
     /// LetOptions is a helper for accepting an arbitrary number of order independent options for
@@ -2055,9 +2065,9 @@
     /// @param name the variable name
     /// @param options the extra options passed to the ast::Var initializer
     /// Can be any of the following, in any order:
-    ///   * ast::Type*          - specifies the variable type
-    ///   * type::AddressSpace   - specifies the variable address space
-    ///   * type::Access         - specifies the variable's access control
+    ///   * ast::Type           - specifies the variable's type
+    ///   * type::AddressSpace  - specifies the variable's address space
+    ///   * type::Access        - specifies the variable's access control
     ///   * ast::Expression*    - specifies the variable's initializer expression
     ///   * ast::Attribute*     - specifies the variable's attributes (repeatable, or vector)
     /// Note that non-repeatable arguments of the same type will use the last argument's value.
@@ -2072,16 +2082,16 @@
     /// @param name the variable name
     /// @param options the extra options passed to the ast::Var initializer
     /// Can be any of the following, in any order:
-    ///   * ast::Type*          - specifies the variable type
-    ///   * type::AddressSpace   - specifies the variable address space
-    ///   * type::Access         - specifies the variable's access control
+    ///   * ast::Type           - specifies the variable's type
+    ///   * type::AddressSpace  - specifies the variable's address space
+    ///   * type::Access        - specifies the variable's access control
     ///   * ast::Expression*    - specifies the variable's initializer expression
     ///   * ast::Attribute*     - specifies the variable's attributes (repeatable, or vector)
     /// Note that non-repeatable arguments of the same type will use the last argument's value.
     /// @returns a `ast::Var` with the given name, address_space and type
     template <typename NAME, typename... OPTIONS>
     const ast::Var* Var(const Source& source, NAME&& name, OPTIONS&&... options) {
-        VarOptions opts(std::forward<OPTIONS>(options)...);
+        VarOptions opts(*this, std::forward<OPTIONS>(options)...);
         return create<ast::Var>(source, Ident(std::forward<NAME>(name)), opts.type,
                                 opts.address_space, opts.access, opts.initializer,
                                 std::move(opts.attributes));
@@ -2091,8 +2101,7 @@
     /// @param options the extra options passed to the ast::Var initializer
     /// Can be any of the following, in any order:
     ///   * ast::Expression*    - specifies the variable's initializer expression (required)
-    ///   * ast::Identifier*    - specifies the variable type
-    ///   * ast::Type*          - specifies the variable type
+    ///   * ast::Type           - specifies the variable's type
     ///   * ast::Attribute*     - specifies the variable's attributes (repeatable, or vector)
     /// Note that non-repeatable arguments of the same type will use the last argument's value.
     /// @returns an `ast::Const` with the given name, type and additional options
@@ -2106,8 +2115,8 @@
     /// @param options the extra options passed to the ast::Var initializer
     /// Can be any of the following, in any order:
     ///   * ast::Expression*    - specifies the variable's initializer expression (required)
-    ///   * ast::Identifier*    - specifies the variable type
-    ///   * ast::Type*          - specifies the variable type
+    ///   * ast::Identifier*    - specifies the variable's type
+    ///   * ast::Type           - specifies the variable's type
     ///   * ast::Attribute*     - specifies the variable's attributes (repeatable, or vector)
     /// Note that non-repeatable arguments of the same type will use the last argument's value.
     /// @returns an `ast::Const` with the given name, type and additional options
@@ -2122,7 +2131,7 @@
     /// @param options the extra options passed to the ast::Var initializer
     /// Can be any of the following, in any order:
     ///   * ast::Expression*    - specifies the variable's initializer expression (required)
-    ///   * ast::Type*          - specifies the variable type
+    ///   * ast::Type           - specifies the variable's type
     ///   * ast::Attribute*     - specifies the variable's attributes (repeatable, or vector)
     /// Note that non-repeatable arguments of the same type will use the last argument's value.
     /// @returns an `ast::Let` with the given name, type and additional options
@@ -2136,7 +2145,7 @@
     /// @param options the extra options passed to the ast::Var initializer
     /// Can be any of the following, in any order:
     ///   * ast::Expression*    - specifies the variable's initializer expression (required)
-    ///   * ast::Type*          - specifies the variable type
+    ///   * ast::Type           - specifies the variable's type
     ///   * ast::Attribute*     - specifies the variable's attributes (repeatable, or vector)
     /// Note that non-repeatable arguments of the same type will use the last argument's value.
     /// @returns an `ast::Let` with the given name, type and additional options
@@ -2175,7 +2184,7 @@
     /// @param name the variable name
     /// @param options the extra options passed to the ast::Var initializer
     /// Can be any of the following, in any order:
-    ///   * ast::Type*          - specifies the variable type
+    ///   * ast::Type           - specifies the variable's type
     ///   * type::AddressSpace   - specifies the variable address space
     ///   * type::Access         - specifies the variable's access control
     ///   * ast::Expression*    - specifies the variable's initializer expression
@@ -2192,7 +2201,7 @@
     /// @param name the variable name
     /// @param options the extra options passed to the ast::Var initializer
     /// Can be any of the following, in any order:
-    ///   * ast::Type*          - specifies the variable type
+    ///   * ast::Type           - specifies the variable's type
     ///   * type::AddressSpace   - specifies the variable address space
     ///   * type::Access         - specifies the variable's access control
     ///   * ast::Expression*    - specifies the variable's initializer expression
@@ -2211,7 +2220,7 @@
     /// @param options the extra options passed to the ast::Const initializer
     /// Can be any of the following, in any order:
     ///   * ast::Expression*    - specifies the variable's initializer expression (required)
-    ///   * ast::Type*          - specifies the variable type
+    ///   * ast::Type           - specifies the variable's type
     ///   * ast::Attribute*     - specifies the variable's attributes (repeatable, or vector)
     /// Note that non-repeatable arguments of the same type will use the last argument's value.
     /// @returns an `ast::Const` with the given name, type and additional options, which is
@@ -2226,7 +2235,7 @@
     /// @param options the extra options passed to the ast::Const initializer
     /// Can be any of the following, in any order:
     ///   * ast::Expression*    - specifies the variable's initializer expression (required)
-    ///   * ast::Type*          - specifies the variable type
+    ///   * ast::Type           - specifies the variable's type
     ///   * ast::Attribute*     - specifies the variable's attributes (repeatable, or vector)
     /// Note that non-repeatable arguments of the same type will use the last argument's value.
     /// @returns an `ast::Const` with the given name, type and additional options, which is
@@ -2242,7 +2251,7 @@
     /// @param options the extra options passed to the ast::Override initializer
     /// Can be any of the following, in any order:
     ///   * ast::Expression*    - specifies the variable's initializer expression (required)
-    ///   * ast::Type*          - specifies the variable type
+    ///   * ast::Type           - specifies the variable's type
     ///   * ast::Attribute*     - specifies the variable's attributes (repeatable, or vector)
     /// Note that non-repeatable arguments of the same type will use the last argument's value.
     /// @returns an `ast::Override` with the given name, type and additional options, which is
@@ -2257,7 +2266,7 @@
     /// @param options the extra options passed to the ast::Override initializer
     /// Can be any of the following, in any order:
     ///   * ast::Expression*    - specifies the variable's initializer expression (required)
-    ///   * ast::Type*          - specifies the variable type
+    ///   * ast::Type           - specifies the variable's type
     ///   * ast::Attribute*     - specifies the variable's attributes (repeatable, or vector)
     /// Note that non-repeatable arguments of the same type will use the last argument's value.
     /// @returns an `ast::Override` with the given name, type and additional options, which is
diff --git a/src/tint/reader/spirv/function.cc b/src/tint/reader/spirv/function.cc
index 046780f..513b7f3 100644
--- a/src/tint/reader/spirv/function.cc
+++ b/src/tint/reader/spirv/function.cc
@@ -2522,7 +2522,8 @@
             }
         }
         auto* var = parser_impl_.MakeVar(inst.result_id(), type::AddressSpace::kUndefined,
-                                         var_store_type, initializer, AttributeList{});
+                                         type::Access::kUndefined, var_store_type, initializer,
+                                         AttributeList{});
         auto* var_decl_stmt = create<ast::VariableDeclStatement>(Source{}, var);
         AddStatement(var_decl_stmt);
         auto* var_type = ty_.Reference(var_store_type, type::AddressSpace::kUndefined);
@@ -3367,8 +3368,9 @@
         // 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, type::AddressSpace::kUndefined, store_type, nullptr,
-                                           AttributeList{})));
+            Source{},
+            parser_impl_.MakeVar(id, type::AddressSpace::kUndefined, type::Access::kUndefined,
+                                 store_type, nullptr, AttributeList{})));
         auto* type = ty_.Reference(store_type, type::AddressSpace::kUndefined);
         identifier_types_.emplace(id, type);
     }
@@ -4835,9 +4837,8 @@
         // either variables or function parameters.
         switch (opcode(inst)) {
             case spv::Op::OpVariable: {
-                if (const auto* module_var = parser_impl_.GetModuleVariable(id)) {
-                    return DefInfo::Pointer{module_var->declared_address_space,
-                                            module_var->declared_access};
+                if (auto v = parser_impl_.GetModuleVariable(id); v.var) {
+                    return DefInfo::Pointer{v.address_space, v.access};
                 }
                 // Local variables are always Function storage class, with default
                 // access mode.
diff --git a/src/tint/reader/spirv/parser_impl.cc b/src/tint/reader/spirv/parser_impl.cc
index 5707bf1..443b28f 100644
--- a/src/tint/reader/spirv/parser_impl.cc
+++ b/src/tint/reader/spirv/parser_impl.cc
@@ -1504,12 +1504,15 @@
             // here.)
             ast_initializer = MakeConstantExpression(var.GetSingleWordInOperand(1)).expr;
         }
-        auto* ast_var = MakeVar(var.result_id(), ast_address_space, ast_store_type, ast_initializer,
-                                utils::Empty);
+        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, utils::Empty);
         // TODO(dneto): initializers (a.k.a. initializer expression)
         if (ast_var) {
             builder_.AST().AddGlobalVariable(ast_var);
-            module_variable_.GetOrCreate(var.result_id(), [ast_var] { return ast_var; });
+            module_variable_.GetOrCreate(var.result_id(), [&] {
+                return ModuleVariable{ast_var, ast_address_space, ast_access};
+            });
         }
     }
 
@@ -1536,14 +1539,16 @@
                                   << init->PrettyPrint();
             }
         }
-        auto* ast_var =
-            MakeVar(builtin_position_.per_vertex_var_id,
-                    enum_converter_.ToAddressSpace(builtin_position_.storage_class),
-                    ConvertType(builtin_position_.position_member_type_id), ast_initializer, {});
+        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,
+                                storage_type, ast_initializer, {});
 
         builder_.AST().AddGlobalVariable(ast_var);
-        module_variable_.GetOrCreate(builtin_position_.per_vertex_var_id,
-                                     [ast_var] { return ast_var; });
+        module_variable_.GetOrCreate(builtin_position_.per_vertex_var_id, [&] {
+            return ModuleVariable{ast_var, ast_address_space};
+        });
     }
     return success_;
 }
@@ -1571,8 +1576,23 @@
     return size->AsIntConstant();
 }
 
+type::Access ParserImpl::VarAccess(const Type* storage_type, type::AddressSpace address_space) {
+    if (address_space != type::AddressSpace::kStorage) {
+        return type::Access::kUndefined;
+    }
+
+    bool read_only = false;
+    if (auto* tn = storage_type->As<Named>()) {
+        read_only = read_only_struct_types_.count(tn->name) > 0;
+    }
+
+    // Apply the access(read) or access(read_write) modifier.
+    return read_only ? type::Access::kRead : type::Access::kReadWrite;
+}
+
 const ast::Var* ParserImpl::MakeVar(uint32_t id,
                                     type::AddressSpace address_space,
+                                    type::Access access,
                                     const Type* storage_type,
                                     const ast::Expression* initializer,
                                     AttributeList decorations) {
@@ -1581,17 +1601,6 @@
         return nullptr;
     }
 
-    type::Access access = type::Access::kUndefined;
-    if (address_space == type::AddressSpace::kStorage) {
-        bool read_only = false;
-        if (auto* tn = storage_type->As<Named>()) {
-            read_only = read_only_struct_types_.count(tn->name) > 0;
-        }
-
-        // Apply the access(read) or access(read_write) modifier.
-        access = read_only ? type::Access::kRead : type::Access::kReadWrite;
-    }
-
     // Handle variables (textures and samplers) are always in the handle
     // address space, so we don't mention the address space.
     if (address_space == type::AddressSpace::kHandle) {
diff --git a/src/tint/reader/spirv/parser_impl.h b/src/tint/reader/spirv/parser_impl.h
index 5d6daa4..0c77a02 100644
--- a/src/tint/reader/spirv/parser_impl.h
+++ b/src/tint/reader/spirv/parser_impl.h
@@ -420,10 +420,17 @@
     /// @returns a list of SPIR-V decorations.
     DecorationList GetMemberPipelineDecorations(const Struct& struct_type, int member_index);
 
+    /// @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.
+    type::Access VarAccess(const Type* storage_type, type::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 type::AddressSpace::kUndefined
+    /// @param access the access
     /// @param storage_type the storage type of the variable
     /// @param initializer the variable initializer
     /// @param decorations the variable decorations
@@ -431,6 +438,7 @@
     /// in the error case
     const ast::Var* MakeVar(uint32_t id,
                             type::AddressSpace address_space,
+                            type::Access access,
                             const Type* storage_type,
                             const ast::Expression* initializer,
                             AttributeList decorations);
@@ -659,13 +667,23 @@
     /// error
     const Type* GetHandleTypeForSpirvHandle(const spvtools::opt::Instruction& obj);
 
+    /// ModuleVariable describes a module scope variable
+    struct ModuleVariable {
+        /// The AST variable node.
+        const ast::Var* var = nullptr;
+        /// The address space of the var
+        type::AddressSpace address_space = type::AddressSpace::kUndefined;
+        /// The access mode of the var
+        type::Access access = type::Access::kUndefined;
+    };
+
     /// Returns the AST variable for the SPIR-V ID of a module-scope variable,
     /// or null if there isn't one.
     /// @param id a SPIR-V ID
     /// @returns the AST variable or null.
-    const ast::Var* GetModuleVariable(uint32_t id) {
+    ModuleVariable GetModuleVariable(uint32_t id) {
         auto entry = module_variable_.Find(id);
-        return entry ? *entry : nullptr;
+        return entry ? *entry : ModuleVariable{};
     }
 
     /// Returns the channel component type corresponding to the given image
@@ -885,7 +903,7 @@
     std::unordered_map<const spvtools::opt::Instruction*, const Type*> handle_type_;
 
     /// Maps the SPIR-V ID of a module-scope variable to its AST variable.
-    utils::Hashmap<uint32_t, const ast::Var*, 16> module_variable_;
+    utils::Hashmap<uint32_t, ModuleVariable, 16> module_variable_;
 
     // Set of symbols of declared type that have been added, used to avoid
     // adding duplicates.
diff --git a/src/tint/reader/wgsl/classify_template_args.cc b/src/tint/reader/wgsl/classify_template_args.cc
index 0ccef95..79d39d2 100644
--- a/src/tint/reader/wgsl/classify_template_args.cc
+++ b/src/tint/reader/wgsl/classify_template_args.cc
@@ -70,8 +70,8 @@
 
     for (size_t i = 0; i < count - 1; i++) {
         switch (tokens[i].type()) {
-            // <identifier> + all type / builtin keywords that will become identifiers.
             case Token::Type::kIdentifier:
+            case Token::Type::kVar:
             case Token::Type::kBitcast: {
                 auto& next = tokens[i + 1];
                 if (next.type() == Token::Type::kLessThan) {
diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index c503353..c9a8576 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -204,23 +204,6 @@
 ParserImpl::FunctionHeader& ParserImpl::FunctionHeader::operator=(const FunctionHeader& rhs) =
     default;
 
-ParserImpl::VarDeclInfo::VarDeclInfo() = default;
-
-ParserImpl::VarDeclInfo::VarDeclInfo(const VarDeclInfo&) = default;
-
-ParserImpl::VarDeclInfo::VarDeclInfo(Source source_in,
-                                     std::string name_in,
-                                     type::AddressSpace address_space_in,
-                                     type::Access access_in,
-                                     ast::Type type_in)
-    : source(std::move(source_in)),
-      name(std::move(name_in)),
-      address_space(address_space_in),
-      access(access_in),
-      type(type_in) {}
-
-ParserImpl::VarDeclInfo::~VarDeclInfo() = default;
-
 ParserImpl::ParserImpl(Source::File const* file) : file_(file) {}
 
 ParserImpl::~ParserImpl() = default;
@@ -614,13 +597,13 @@
 
     TINT_DEFER(attrs.Clear());
 
-    return builder_.Var(decl->source,         // source
-                        decl->name,           // symbol
-                        decl->type,           // type
-                        decl->address_space,  // address space
-                        decl->access,         // access control
-                        initializer,          // initializer
-                        std::move(attrs));    // attributes
+    return builder_.create<ast::Var>(decl->source,                // source
+                                     builder_.Ident(decl->name),  // symbol
+                                     decl->type,                  // type
+                                     decl->address_space,         // address space
+                                     decl->access,                // access control
+                                     initializer,                 // initializer
+                                     std::move(attrs));           // attributes
 }
 
 // global_constant_decl :
@@ -756,37 +739,28 @@
     return expect_ident_with_optional_type_specifier(use, /* allow_inferred */ false);
 }
 
-// access_mode
-//   : 'read'
-//   | 'write'
-//   | 'read_write'
-Expect<type::Access> ParserImpl::expect_access_mode(std::string_view use) {
-    return expect_enum("access control", type::ParseAccess, type::kAccessStrings, use);
-}
-
 // variable_qualifier
-//   : LESS_THAN address_spaces (COMMA access_mode)? GREATER_THAN
+//   : _template_args_start expression (COMMA expression)? _template_args_end
 Maybe<ParserImpl::VariableQualifier> ParserImpl::variable_qualifier() {
-    if (!peek_is(Token::Type::kLessThan)) {
+    if (!peek_is(Token::Type::kTemplateArgsLeft) && !peek_is(Token::Type::kLessThan)) {
+        // Note: kLessThan will give a sensible error at expect_template_arg_block()
         return Failure::kNoMatch;
     }
 
     auto* use = "variable declaration";
-    auto vq = expect_lt_gt_block(use, [&]() -> Expect<VariableQualifier> {
-        auto source = make_source_range();
-        auto sc = expect_address_space(use);
-        if (sc.errored) {
+    auto vq = expect_template_arg_block(use, [&]() -> Expect<VariableQualifier> {
+        auto address_space = expect_expression("'var' address space");
+        if (address_space.errored) {
             return Failure::kErrored;
         }
         if (match(Token::Type::kComma)) {
-            auto ac = expect_access_mode(use);
-            if (ac.errored) {
+            auto access = expect_expression("'var' access mode");
+            if (access.errored) {
                 return Failure::kErrored;
             }
-            return VariableQualifier{sc.value, ac.value};
+            return VariableQualifier{address_space.value, access.value};
         }
-        return Expect<VariableQualifier>{VariableQualifier{sc.value, type::Access::kUndefined},
-                                         source};
+        return VariableQualifier{address_space.value};
     });
 
     if (vq.errored) {
@@ -900,18 +874,6 @@
     return type.value;
 }
 
-// address_space
-//   : 'function'
-//   | 'private'
-//   | 'workgroup'
-//   | 'uniform'
-//   | 'storage'
-//
-// Note, we also parse `push_constant` from the experimental extension
-Expect<type::AddressSpace> ParserImpl::expect_address_space(std::string_view use) {
-    return expect_enum("address space", type::ParseAddressSpace, type::kAddressSpaceStrings, use);
-}
-
 // struct_decl
 //   : STRUCT IDENT struct_body_decl
 Maybe<const ast::Struct*> ParserImpl::struct_decl() {
@@ -1519,12 +1481,13 @@
         initializer = initializer_expr.value;
     }
 
-    auto* var = builder_.Var(decl_source,          // source
-                             decl->name,           // symbol
-                             decl->type,           // type
-                             decl->address_space,  // address space
-                             decl->access,         // access control
-                             initializer);         // initializer
+    auto* var = builder_.create<ast::Var>(decl_source,                 // source
+                                          builder_.Ident(decl->name),  // symbol
+                                          decl->type,                  // type
+                                          decl->address_space,         // address space
+                                          decl->access,                // access control
+                                          initializer,                 // initializer
+                                          utils::Empty);               // attributes
 
     return create<ast::VariableDeclStatement>(var->source, var);
 }
@@ -2520,7 +2483,7 @@
     return create<ast::BinaryExpression>(tok_op.source(), op, lhs, rhs.value);
 }
 
-Expect<const ast::Expression*> ParserImpl::expect_expression() {
+Expect<const ast::Expression*> ParserImpl::expect_expression(std::string_view use) {
     auto& t = peek();
     auto expr = expression();
     if (expr.errored) {
@@ -2529,7 +2492,7 @@
     if (expr.matched) {
         return expr.value;
     }
-    return add_error(t, "expected expression");
+    return add_error(t, "expected expression for " + std::string(use));
 }
 
 Expect<utils::Vector<const ast::Expression*, 3>> ParserImpl::expect_expression_list(
@@ -2537,7 +2500,7 @@
     Token::Type terminator) {
     utils::Vector<const ast::Expression*, 3> exprs;
     while (continue_parsing()) {
-        auto expr = expect_expression();
+        auto expr = expect_expression(use);
         if (expr.errored) {
             return Failure::kErrored;
         }
diff --git a/src/tint/reader/wgsl/parser_impl.h b/src/tint/reader/wgsl/parser_impl.h
index f692d77..64618b3 100644
--- a/src/tint/reader/wgsl/parser_impl.h
+++ b/src/tint/reader/wgsl/parser_impl.h
@@ -148,10 +148,9 @@
 
         /// Constructor for a successful parse.
         /// @param val the result value of the parse
-        /// @param s the optional source of the value
         template <typename U>
-        inline Maybe(U&& val, const Source& s = {})  // NOLINT
-            : value(std::forward<U>(val)), source(s), matched(true) {}
+        inline Maybe(U&& val)  // NOLINT
+            : value(std::forward<U>(val)), matched(true) {}
 
         /// Constructor for parse error state.
         inline Maybe(Failure::Errored) : errored(true) {}  // NOLINT
@@ -163,16 +162,13 @@
         /// @param e the Expect to copy this Maybe from
         template <typename U>
         inline Maybe(const Expect<U>& e)  // NOLINT
-            : value(e.value), source(e.value), errored(e.errored), matched(!e.errored) {}
+            : value(e.value), errored(e.errored), matched(!e.errored) {}
 
         /// Move from an Expect.
         /// @param e the Expect to move this Maybe from
         template <typename U>
         inline Maybe(Expect<U>&& e)  // NOLINT
-            : value(std::move(e.value)),
-              source(std::move(e.source)),
-              errored(e.errored),
-              matched(!e.errored) {}
+            : value(std::move(e.value)), errored(e.errored), matched(!e.errored) {}
 
         /// Copy constructor
         inline Maybe(const Maybe&) = default;
@@ -197,8 +193,6 @@
         /// The value of a successful parse.
         /// Zero-initialized when there was a parse error.
         T value{};
-        /// Optional source of the value.
-        Source source;
         /// True if there was a error parsing.
         bool errored = false;
         /// True if there was a error parsing.
@@ -268,33 +262,14 @@
 
     /// VarDeclInfo contains the parsed information for variable declaration.
     struct VarDeclInfo {
-        /// Constructor
-        VarDeclInfo();
-        /// Copy constructor
-        /// @param other the VarDeclInfo to copy
-        VarDeclInfo(const VarDeclInfo& other);
-        /// Constructor
-        /// @param source_in variable declaration source
-        /// @param name_in variable name
-        /// @param address_space_in variable address space
-        /// @param access_in variable access control
-        /// @param type_in variable type
-        VarDeclInfo(Source source_in,
-                    std::string name_in,
-                    type::AddressSpace address_space_in,
-                    type::Access access_in,
-                    ast::Type type_in);
-        /// Destructor
-        ~VarDeclInfo();
-
         /// Variable declaration source
         Source source;
         /// Variable name
         std::string name;
         /// Variable address space
-        type::AddressSpace address_space = type::AddressSpace::kUndefined;
+        const ast::Expression* address_space = nullptr;
         /// Variable access control
-        type::Access access = type::Access::kUndefined;
+        const ast::Expression* access = nullptr;
         /// Variable type
         ast::Type type;
     };
@@ -302,9 +277,9 @@
     /// VariableQualifier contains the parsed information for a variable qualifier
     struct VariableQualifier {
         /// The variable's address space
-        type::AddressSpace address_space = type::AddressSpace::kUndefined;
+        const ast::Expression* address_space = nullptr;
         /// The variable's access control
-        type::Access access = type::Access::kUndefined;
+        const ast::Expression* access = nullptr;
     };
 
     /// MatrixDimensions contains the column and row information for a matrix
@@ -447,10 +422,6 @@
     /// Parses a `type_specifier` grammar element
     /// @returns the parsed Type or nullptr if none matched.
     Maybe<ast::Type> type_specifier();
-    /// Parses an `address_space` grammar element, erroring on parse failure.
-    /// @param use a description of what was being parsed if an error was raised.
-    /// @returns the address space or type::AddressSpace::kUndefined if none matched
-    Expect<type::AddressSpace> expect_address_space(std::string_view use);
     /// Parses a `struct_decl` grammar element.
     /// @returns the struct type or nullptr on error
     Maybe<const ast::Struct*> struct_decl();
@@ -482,11 +453,6 @@
     /// not match a stage name.
     /// @returns the pipeline stage.
     Expect<ast::PipelineStage> expect_pipeline_stage();
-    /// Parses an access control identifier, erroring if the next token does not
-    /// match a valid access control.
-    /// @param use a description of what was being parsed if an error was raised
-    /// @returns the parsed access control.
-    Expect<type::Access> expect_access_mode(std::string_view use);
     /// Parses an interpolation sample name identifier, erroring if the next token does not match a
     /// valid sample name.
     /// @returns the parsed sample name.
@@ -597,8 +563,9 @@
     /// @returns the parsed expression or nullptr
     Maybe<const ast::Expression*> expression();
     /// Parses the `expression` grammar rule
+    /// @param use the use of the expression
     /// @returns the parsed expression or error
-    Expect<const ast::Expression*> expect_expression();
+    Expect<const ast::Expression*> expect_expression(std::string_view use);
     /// Parses a comma separated expression list
     /// @param use the use of the expression list
     /// @param terminator the terminating token for the list
diff --git a/src/tint/reader/wgsl/parser_impl_address_space_test.cc b/src/tint/reader/wgsl/parser_impl_address_space_test.cc
deleted file mode 100644
index 32c3c47..0000000
--- a/src/tint/reader/wgsl/parser_impl_address_space_test.cc
+++ /dev/null
@@ -1,62 +0,0 @@
-// Copyright 2020 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "src/tint/reader/wgsl/parser_impl_test_helper.h"
-
-namespace tint::reader::wgsl {
-namespace {
-
-struct AddressSpaceData {
-    const char* input;
-    type::AddressSpace result;
-};
-inline std::ostream& operator<<(std::ostream& out, AddressSpaceData data) {
-    out << std::string(data.input);
-    return out;
-}
-
-class ParserAddressSpaceTest : public ParserImplTestWithParam<AddressSpaceData> {};
-
-TEST_P(ParserAddressSpaceTest, Parses) {
-    auto params = GetParam();
-    auto p = parser(params.input);
-
-    auto sc = p->expect_address_space("test");
-    EXPECT_FALSE(sc.errored);
-    EXPECT_FALSE(p->has_error());
-    EXPECT_EQ(sc.value, params.result);
-
-    auto& t = p->next();
-    EXPECT_TRUE(t.IsEof());
-}
-INSTANTIATE_TEST_SUITE_P(
-    ParserImplTest,
-    ParserAddressSpaceTest,
-    testing::Values(AddressSpaceData{"uniform", type::AddressSpace::kUniform},
-                    AddressSpaceData{"workgroup", type::AddressSpace::kWorkgroup},
-                    AddressSpaceData{"storage", type::AddressSpace::kStorage},
-                    AddressSpaceData{"private", type::AddressSpace::kPrivate},
-                    AddressSpaceData{"function", type::AddressSpace::kFunction}));
-
-TEST_F(ParserImplTest, AddressSpace_NoMatch) {
-    auto p = parser("not-a-address-space");
-    auto sc = p->expect_address_space("test");
-    EXPECT_EQ(sc.errored, true);
-    EXPECT_TRUE(p->has_error());
-    EXPECT_EQ(p->error(), R"(1:1: expected address space for test
-Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup')");
-}
-
-}  // namespace
-}  // namespace tint::reader::wgsl
diff --git a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
index d40898a..18b40be 100644
--- a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
@@ -799,7 +799,7 @@
 
 TEST_F(ParserImplErrorTest, GlobalDeclStorageTextureMissingSubtype) {
     EXPECT("var x : texture_storage_2d<>;",
-           R"(test.wgsl:1:28 error: expected expression
+           R"(test.wgsl:1:28 error: expected expression for type template argument list
 var x : texture_storage_2d<>;
                            ^
 )");
@@ -1123,20 +1123,11 @@
 )");
 }
 
-TEST_F(ParserImplErrorTest, GlobalDeclVarStorageDeclInvalidClass) {
-    EXPECT("var<fish> i : i32",
-           R"(test.wgsl:1:5 error: expected address space for variable declaration
-Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup'
-var<fish> i : i32
-    ^^^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclVarStorageDeclMissingGThan) {
     EXPECT("var<private i : i32",
-           R"(test.wgsl:1:13 error: expected '>' for variable declaration
+           R"(test.wgsl:1:4 error: missing closing '>' for variable declaration
 var<private i : i32
-            ^
+   ^
 )");
 }
 
diff --git a/src/tint/reader/wgsl/parser_impl_global_variable_decl_test.cc b/src/tint/reader/wgsl/parser_impl_global_variable_decl_test.cc
index 92272cc..daf978e 100644
--- a/src/tint/reader/wgsl/parser_impl_global_variable_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_global_variable_decl_test.cc
@@ -31,10 +31,8 @@
     ASSERT_NE(var, nullptr);
 
     ast::CheckIdentifier(p->builder().Symbols(), var->name, "a");
-
     ast::CheckIdentifier(p->builder().Symbols(), var->type, "f32");
-
-    EXPECT_EQ(var->declared_address_space, type::AddressSpace::kPrivate);
+    ast::CheckIdentifier(p->builder().Symbols(), var->declared_address_space, "private");
 
     EXPECT_EQ(var->source.range.begin.line, 1u);
     EXPECT_EQ(var->source.range.begin.column, 14u);
@@ -58,8 +56,7 @@
 
     ast::CheckIdentifier(p->builder().Symbols(), var->name, "a");
     ast::CheckIdentifier(p->builder().Symbols(), var->type, "f32");
-
-    EXPECT_EQ(var->declared_address_space, type::AddressSpace::kPrivate);
+    ast::CheckIdentifier(p->builder().Symbols(), var->declared_address_space, "private");
 
     EXPECT_EQ(var->source.range.begin.line, 1u);
     EXPECT_EQ(var->source.range.begin.column, 14u);
@@ -83,11 +80,8 @@
     ASSERT_NE(var, nullptr);
 
     ast::CheckIdentifier(p->builder().Symbols(), var->name, "a");
-    ASSERT_NE(var->type, nullptr);
-
     ast::CheckIdentifier(p->builder().Symbols(), var->type, "f32");
-
-    EXPECT_EQ(var->declared_address_space, type::AddressSpace::kUniform);
+    ast::CheckIdentifier(p->builder().Symbols(), var->declared_address_space, "uniform");
 
     EXPECT_EQ(var->source.range.begin.line, 1u);
     EXPECT_EQ(var->source.range.begin.column, 36u);
@@ -116,10 +110,8 @@
     ASSERT_NE(var, nullptr);
 
     ast::CheckIdentifier(p->builder().Symbols(), var->name, "a");
-    ASSERT_NE(var->type, nullptr);
     ast::CheckIdentifier(p->builder().Symbols(), var->type, "f32");
-
-    EXPECT_EQ(var->declared_address_space, type::AddressSpace::kUniform);
+    ast::CheckIdentifier(p->builder().Symbols(), var->declared_address_space, "uniform");
 
     EXPECT_EQ(var->source.range.begin.line, 1u);
     EXPECT_EQ(var->source.range.begin.column, 36u);
@@ -162,19 +154,5 @@
     EXPECT_EQ(p->error(), "1:24: missing initializer for 'var' declaration");
 }
 
-TEST_F(ParserImplTest, GlobalVariableDecl_InvalidVariableDecl) {
-    auto p = parser("var<invalid> a : f32;");
-    auto attrs = p->attribute_list();
-    EXPECT_FALSE(attrs.errored);
-    EXPECT_FALSE(attrs.matched);
-    auto e = p->global_variable_decl(attrs.value);
-    EXPECT_TRUE(p->has_error());
-    EXPECT_TRUE(e.errored);
-    EXPECT_FALSE(e.matched);
-    EXPECT_EQ(e.value, nullptr);
-    EXPECT_EQ(p->error(), R"(1:5: expected address space for variable declaration
-Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup')");
-}
-
 }  // namespace
 }  // namespace tint::reader::wgsl
diff --git a/src/tint/reader/wgsl/parser_impl_type_decl_test.cc b/src/tint/reader/wgsl/parser_impl_type_decl_test.cc
index b40d887..25e60e0 100644
--- a/src/tint/reader/wgsl/parser_impl_type_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_type_decl_test.cc
@@ -137,7 +137,7 @@
     EXPECT_FALSE(t.matched);
     ASSERT_EQ(t.value, nullptr);
     ASSERT_TRUE(p->has_error());
-    ASSERT_EQ(p->error(), "1:6: expected expression");
+    ASSERT_EQ(p->error(), "1:6: expected expression for type template argument list");
 }
 INSTANTIATE_TEST_SUITE_P(ParserImplTest,
                          VecMissingType,
@@ -211,7 +211,7 @@
     EXPECT_FALSE(t.matched);
     ASSERT_EQ(t.value, nullptr);
     ASSERT_TRUE(p->has_error());
-    ASSERT_EQ(p->error(), R"(1:5: expected expression)");
+    ASSERT_EQ(p->error(), R"(1:5: expected expression for type template argument list)");
 }
 
 TEST_F(ParserImplTest, TypeDecl_Ptr_MissingParams) {
@@ -221,7 +221,7 @@
     EXPECT_FALSE(t.matched);
     ASSERT_EQ(t.value, nullptr);
     ASSERT_TRUE(p->has_error());
-    ASSERT_EQ(p->error(), R"(1:5: expected expression)");
+    ASSERT_EQ(p->error(), R"(1:5: expected expression for type template argument list)");
 }
 
 TEST_F(ParserImplTest, TypeDecl_Atomic) {
@@ -256,7 +256,7 @@
     EXPECT_FALSE(t.matched);
     ASSERT_EQ(t.value, nullptr);
     ASSERT_TRUE(p->has_error());
-    ASSERT_EQ(p->error(), "1:8: expected expression");
+    ASSERT_EQ(p->error(), "1:8: expected expression for type template argument list");
 }
 
 TEST_F(ParserImplTest, TypeDecl_Array_AbstractIntLiteralSize) {
@@ -431,7 +431,7 @@
     EXPECT_FALSE(t.matched);
     ASSERT_EQ(t.value, nullptr);
     ASSERT_TRUE(p->has_error());
-    ASSERT_EQ(p->error(), "1:8: expected expression");
+    ASSERT_EQ(p->error(), "1:8: expected expression for type template argument list");
 }
 INSTANTIATE_TEST_SUITE_P(ParserImplTest,
                          MatrixMissingType,
diff --git a/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc b/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
index 9138d55..993ecb0 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
@@ -84,8 +84,7 @@
     EXPECT_EQ(v->name, "my_var");
 
     ast::CheckIdentifier(p->builder().Symbols(), v->type, "f32");
-
-    EXPECT_EQ(v->address_space, type::AddressSpace::kPrivate);
+    ast::CheckIdentifier(p->builder().Symbols(), v->address_space, "private");
 
     EXPECT_EQ(v->source.range.begin.line, 1u);
     EXPECT_EQ(v->source.range.begin.column, 14u);
@@ -102,20 +101,7 @@
     EXPECT_EQ(v->name, "my_var");
 
     ast::CheckIdentifier(p->builder().Symbols(), v->type, "f32");
-
-    EXPECT_EQ(v->address_space, type::AddressSpace::kPushConstant);
-}
-
-TEST_F(ParserImplTest, VariableDecl_InvalidAddressSpace) {
-    auto p = parser("var<unknown> my_var : f32");
-    auto v = p->variable_decl();
-    EXPECT_FALSE(v.matched);
-    EXPECT_TRUE(v.errored);
-    EXPECT_TRUE(p->has_error());
-    EXPECT_EQ(p->error(),
-              R"(1:5: expected address space for variable declaration
-Did you mean 'uniform'?
-Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup')");
+    ast::CheckIdentifier(p->builder().Symbols(), v->address_space, "push_constant");
 }
 
 }  // namespace
diff --git a/src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc b/src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc
index 111f207..9a8b936 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_qualifier_test.cc
@@ -12,6 +12,7 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
+#include "src/tint/ast/test_helper.h"
 #include "src/tint/reader/wgsl/parser_impl_test_helper.h"
 
 namespace tint::reader::wgsl {
@@ -31,14 +32,23 @@
 
 TEST_P(VariableQualifierTest, ParsesAddressSpace) {
     auto params = GetParam();
-    auto p = parser(std::string("<") + params.input + ">");
+    auto p = parser(std::string("var<") + params.input + "> name");
 
-    auto sc = p->variable_qualifier();
+    auto sc = p->variable_decl();
     EXPECT_FALSE(p->has_error());
     EXPECT_FALSE(sc.errored);
     EXPECT_TRUE(sc.matched);
-    EXPECT_EQ(sc->address_space, params.address_space);
-    EXPECT_EQ(sc->access, params.access);
+    if (params.address_space != type::AddressSpace::kUndefined) {
+        ast::CheckIdentifier(p->builder().Symbols(), sc->address_space,
+                             utils::ToString(params.address_space));
+    } else {
+        EXPECT_EQ(sc->address_space, nullptr);
+    }
+    if (params.access != type::Access::kUndefined) {
+        ast::CheckIdentifier(p->builder().Symbols(), sc->access, utils::ToString(params.access));
+    } else {
+        EXPECT_EQ(sc->access, nullptr);
+    }
 
     auto& t = p->next();
     EXPECT_TRUE(t.IsEof());
@@ -57,24 +67,13 @@
         VariableStorageData{"storage, read_write", type::AddressSpace::kStorage,
                             type::Access::kReadWrite}));
 
-TEST_F(ParserImplTest, VariableQualifier_NoMatch) {
-    auto p = parser("<not-a-storage-class>");
-    auto sc = p->variable_qualifier();
-    EXPECT_TRUE(p->has_error());
-    EXPECT_TRUE(sc.errored);
-    EXPECT_FALSE(sc.matched);
-    EXPECT_EQ(p->error(), R"(1:2: expected address space for variable declaration
-Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup')");
-}
-
 TEST_F(ParserImplTest, VariableQualifier_Empty) {
-    auto p = parser("<>");
-    auto sc = p->variable_qualifier();
+    auto p = parser("var<> name");
+    auto sc = p->variable_decl();
     EXPECT_TRUE(p->has_error());
     EXPECT_TRUE(sc.errored);
     EXPECT_FALSE(sc.matched);
-    EXPECT_EQ(p->error(), R"(1:2: expected address space for variable declaration
-Possible values: 'function', 'private', 'push_constant', 'storage', 'uniform', 'workgroup')");
+    EXPECT_EQ(p->error(), R"(1:5: expected expression for 'var' address space)");
 }
 
 TEST_F(ParserImplTest, VariableQualifier_MissingLessThan) {
@@ -105,7 +104,7 @@
     EXPECT_TRUE(p->has_error());
     EXPECT_TRUE(sc.errored);
     EXPECT_FALSE(sc.matched);
-    EXPECT_EQ(p->error(), "1:9: expected '>' for variable declaration");
+    EXPECT_EQ(p->error(), "1:1: missing closing '>' for variable declaration");
 }
 
 }  // namespace
diff --git a/src/tint/resolver/address_space_validation_test.cc b/src/tint/resolver/address_space_validation_test.cc
index 2c879de..95301c7 100644
--- a/src/tint/resolver/address_space_validation_test.cc
+++ b/src/tint/resolver/address_space_validation_test.cc
@@ -32,8 +32,9 @@
     GlobalVar(Source{{12, 34}}, "g", ty.f32());
 
     EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              "12:34 error: module-scope 'var' declaration must have a address space");
+    EXPECT_EQ(
+        r()->error(),
+        R"(12:34 error: module-scope 'var' declarations that are not of texture or sampler types must provide an address space)");
 }
 
 TEST_F(ResolverAddressSpaceValidationTest, PointerAlias_NoAddressSpace_Fail) {
@@ -469,7 +470,7 @@
 
     EXPECT_EQ(
         r()->error(),
-        R"(12:34 error: only variables in <storage> address space may declare an access mode)");
+        R"(12:34 error: only variables in <storage> address space may specify an access mode)");
 }
 
 TEST_F(ResolverAddressSpaceValidationTest, PointerAlias_NotStorage_AccessMode) {
@@ -481,7 +482,7 @@
 
     EXPECT_EQ(
         r()->error(),
-        R"(12:34 error: only pointers in <storage> address space may declare an access mode)");
+        R"(12:34 error: only pointers in <storage> address space may specify an access mode)");
 }
 
 TEST_F(ResolverAddressSpaceValidationTest, GlobalVariable_Storage_ReadAccessMode) {
diff --git a/src/tint/resolver/dependency_graph.cc b/src/tint/resolver/dependency_graph.cc
index 16a3f8e..75e377f 100644
--- a/src/tint/resolver/dependency_graph.cc
+++ b/src/tint/resolver/dependency_graph.cc
@@ -183,13 +183,9 @@
                 Declare(func->name->symbol, func);
                 TraverseFunction(func);
             },
-            [&](const ast::Variable* var) {
-                Declare(var->name->symbol, var);
-                TraverseTypeExpression(var->type);
-                TraverseAttributes(var->attributes);
-                if (var->initializer) {
-                    TraverseValueExpression(var->initializer);
-                }
+            [&](const ast::Variable* v) {
+                Declare(v->name->symbol, v);
+                TraverseVariable(v);
             },
             [&](const ast::DiagnosticDirective*) {
                 // Diagnostic directives do not affect the dependency graph.
@@ -204,8 +200,18 @@
     }
 
   private:
-    /// Traverses the function, performing symbol resolution and determining
-    /// global dependencies.
+    /// Traverses the variable, performing symbol resolution.
+    void TraverseVariable(const ast::Variable* v) {
+        if (auto* var = v->As<ast::Var>()) {
+            TraverseAddressSpaceExpression(var->declared_address_space);
+            TraverseAccessExpression(var->declared_access);
+        }
+        TraverseTypeExpression(v->type);
+        TraverseAttributes(v->attributes);
+        TraverseValueExpression(v->initializer);
+    }
+
+    /// Traverses the function, performing symbol resolution and determining global dependencies.
     void TraverseFunction(const ast::Function* func) {
         TraverseAttributes(func->attributes);
         TraverseAttributes(func->return_type_attributes);
@@ -301,8 +307,7 @@
                 if (auto* shadows = scope_stack_.Get(v->variable->name->symbol)) {
                     graph_.shadows.Add(v->variable, shadows);
                 }
-                TraverseTypeExpression(v->variable->type);
-                TraverseValueExpression(v->variable->initializer);
+                TraverseVariable(v->variable);
                 Declare(v->variable->name->symbol, v->variable);
             },
             [&](const ast::WhileStatement* w) {
@@ -345,6 +350,18 @@
         TraverseExpression(root, "type", "references");
     }
 
+    /// Traverses the expression @p root_expr for the intended use as an address space, performing
+    /// symbol resolution and determining global dependencies.
+    void TraverseAddressSpaceExpression(const ast::Expression* root) {
+        TraverseExpression(root, "address space", "references");
+    }
+
+    /// Traverses the expression @p root_expr for the intended use as an access, performing symbol
+    /// resolution and determining global dependencies.
+    void TraverseAccessExpression(const ast::Expression* root) {
+        TraverseExpression(root, "access", "references");
+    }
+
     /// Traverses the expression @p root_expr for the intended use as a call target, performing
     /// symbol resolution and determining global dependencies.
     void TraverseCallableExpression(const ast::Expression* root) {
diff --git a/src/tint/resolver/dependency_graph_test.cc b/src/tint/resolver/dependency_graph_test.cc
index 31e484e..a316ad0 100644
--- a/src/tint/resolver/dependency_graph_test.cc
+++ b/src/tint/resolver/dependency_graph_test.cc
@@ -1437,13 +1437,13 @@
         << resolved->String(Symbols(), Diagnostics());
 }
 
-TEST_P(ResolverDependencyGraphResolveToAddressSpace, ShadowedByGlobalVar) {
+TEST_P(ResolverDependencyGraphResolveToAddressSpace, ShadowedByGlobalConst) {
     const auto use = std::get<0>(GetParam());
     const auto builtin = std::get<1>(GetParam());
     const auto symbol = Symbols().New(utils::ToString(builtin));
 
     SymbolTestHelper helper(this);
-    auto* decl = helper.Add(SymbolDeclKind::GlobalVar, symbol);
+    auto* decl = helper.Add(SymbolDeclKind::GlobalConst, symbol);
     auto* ident = helper.Add(use, symbol);
     helper.Build();
 
diff --git a/src/tint/resolver/function_validation_test.cc b/src/tint/resolver/function_validation_test.cc
index b04831e..1062958 100644
--- a/src/tint/resolver/function_validation_test.cc
+++ b/src/tint/resolver/function_validation_test.cc
@@ -1088,15 +1088,14 @@
         param.expectation == Expectation::kPassWithFullPtrParameterExtension) {
         ASSERT_TRUE(r()->Resolve()) << r()->error();
     } else {
-        std::stringstream ss;
-        ss << param.address_space;
         EXPECT_FALSE(r()->Resolve());
         if (param.expectation == Expectation::kInvalid) {
-            EXPECT_EQ(r()->error(), "12:34 error: unknown identifier: '" + ss.str() + "'");
+            EXPECT_EQ(r()->error(), "12:34 error: unknown identifier: '" +
+                                        utils::ToString(param.address_space) + "'");
         } else {
             EXPECT_EQ(r()->error(),
-                      "12:34 error: function parameter of pointer type cannot be in '" + ss.str() +
-                          "' address space");
+                      "12:34 error: function parameter of pointer type cannot be in '" +
+                          utils::ToString(param.address_space) + "' address space");
         }
     }
 }
@@ -1105,8 +1104,8 @@
     ResolverFunctionParameterValidationTest,
     testing::Values(
         TestParams{type::AddressSpace::kUndefined, Expectation::kInvalid},
-        TestParams{type::AddressSpace::kIn, Expectation::kInvalid},
-        TestParams{type::AddressSpace::kOut, Expectation::kInvalid},
+        TestParams{type::AddressSpace::kIn, Expectation::kAlwaysFail},
+        TestParams{type::AddressSpace::kOut, Expectation::kAlwaysFail},
         TestParams{type::AddressSpace::kUniform, Expectation::kPassWithFullPtrParameterExtension},
         TestParams{type::AddressSpace::kWorkgroup, Expectation::kPassWithFullPtrParameterExtension},
         TestParams{type::AddressSpace::kHandle, Expectation::kInvalid},
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index b87852f..8336765 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -473,8 +473,14 @@
         return nullptr;
     }
 
-    auto address_space = var->declared_address_space;
-    if (address_space == type::AddressSpace::kUndefined) {
+    auto address_space = type::AddressSpace::kUndefined;
+    if (var->declared_address_space) {
+        auto expr = AddressSpaceExpression(var->declared_address_space);
+        if (!expr) {
+            return nullptr;
+        }
+        address_space = expr->Value();
+    } else {
         // No declared address space. Infer from usage / type.
         if (!is_global) {
             address_space = type::AddressSpace::kFunction;
@@ -494,8 +500,14 @@
         return nullptr;
     }
 
-    auto access = var->declared_access;
-    if (access == type::Access::kUndefined) {
+    auto access = type::Access::kUndefined;
+    if (var->declared_access) {
+        auto expr = AccessExpression(var->declared_access);
+        if (!expr) {
+            return nullptr;
+        }
+        access = expr->Value();
+    } else {
         access = DefaultAccessForAddressSpace(address_space);
     }
 
diff --git a/src/tint/resolver/validation_test.cc b/src/tint/resolver/validation_test.cc
index 73fa9b8..8ee795c 100644
--- a/src/tint/resolver/validation_test.cc
+++ b/src/tint/resolver/validation_test.cc
@@ -295,22 +295,23 @@
 
 TEST_F(ResolverValidationTest, AddressSpace_SamplerExplicitAddressSpace) {
     auto t = ty.sampler(type::SamplerKind::kSampler);
-    GlobalVar(Source{{12, 34}}, "var", t, type::AddressSpace::kHandle, Binding(0_a), Group(0_a));
+    GlobalVar(Source{{12, 34}}, "var", t, type::AddressSpace::kPrivate, Binding(0_a), Group(0_a));
 
     EXPECT_FALSE(r()->Resolve());
 
     EXPECT_EQ(r()->error(),
-              R"(12:34 error: variables of type 'sampler' must not have a address space)");
+              R"(12:34 error: variables of type 'sampler' must not specifiy an address space)");
 }
 
 TEST_F(ResolverValidationTest, AddressSpace_TextureExplicitAddressSpace) {
     auto t = ty.sampled_texture(type::TextureDimension::k1d, ty.f32());
-    GlobalVar(Source{{12, 34}}, "var", t, type::AddressSpace::kHandle, Binding(0_a), Group(0_a));
+    GlobalVar(Source{{12, 34}}, "var", t, type::AddressSpace::kFunction, Binding(0_a), Group(0_a));
 
     EXPECT_FALSE(r()->Resolve()) << r()->error();
 
-    EXPECT_EQ(r()->error(),
-              R"(12:34 error: variables of type 'texture_1d<f32>' must not have a address space)");
+    EXPECT_EQ(
+        r()->error(),
+        R"(12:34 error: variables of type 'texture_1d<f32>' must not specifiy an address space)");
 }
 
 TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadChar) {
diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc
index 5b399cf..2935c70 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -299,7 +299,7 @@
         // * For the storage address space, the access mode is optional, and defaults to read.
         // * For other address spaces, the access mode must not be written.
         if (s->AddressSpace() != type::AddressSpace::kStorage) {
-            AddError("only pointers in <storage> address space may declare an access mode",
+            AddError("only pointers in <storage> address space may specify an access mode",
                      a->source);
             return false;
         }
@@ -604,7 +604,7 @@
     }
     bool ok = Switch(
         decl,  //
-        [&](const ast::Var*) {
+        [&](const ast::Var* var) {
             if (auto* init = global->Initializer();
                 init && init->Stage() > sem::EvaluationStage::kOverride) {
                 AddError("module-scope 'var' initializer must be a constant or override-expression",
@@ -612,8 +612,11 @@
                 return false;
             }
 
-            if (global->AddressSpace() == type::AddressSpace::kUndefined) {
-                AddError("module-scope 'var' declaration must have a address space", decl->source);
+            if (!var->declared_address_space && !global->Type()->UnwrapRef()->is_handle()) {
+                AddError(
+                    "module-scope 'var' declarations that are not of texture or sampler types must "
+                    "provide an address space",
+                    decl->source);
                 return false;
             }
 
@@ -696,25 +699,23 @@
         return false;
     }
 
-    if (store_ty->is_handle()) {
-        if (var->declared_address_space != type::AddressSpace::kUndefined) {
-            // https://gpuweb.github.io/gpuweb/wgsl/#module-scope-variables
-            // If the store type is a texture type or a sampler type, then the variable declaration
-            // must not have a address space attribute. The address space will always be handle.
-            AddError("variables of type '" + sem_.TypeNameOf(store_ty) +
-                         "' must not have a address space",
-                     var->source);
-            return false;
-        }
+    if (store_ty->is_handle() && var->declared_address_space) {
+        // https://gpuweb.github.io/gpuweb/wgsl/#module-scope-variables
+        // If the store type is a texture type or a sampler type, then the variable declaration must
+        // not have a address space attribute. The address space will always be handle.
+        AddError("variables of type '" + sem_.TypeNameOf(store_ty) +
+                     "' must not specifiy an address space",
+                 var->source);
+        return false;
     }
 
-    if (var->declared_access != type::Access::kUndefined) {
+    if (var->declared_access) {
         // https://www.w3.org/TR/WGSL/#access-mode-defaults
         // When writing a variable declaration or a pointer type in WGSL source:
         // * For the storage address space, the access mode is optional, and defaults to read.
         // * For other address spaces, the access mode must not be written.
-        if (var->declared_address_space != type::AddressSpace::kStorage) {
-            AddError("only variables in <storage> address space may declare an access mode",
+        if (v->AddressSpace() != type::AddressSpace::kStorage) {
+            AddError("only variables in <storage> address space may specify an access mode",
                      var->source);
             return false;
         }
@@ -726,8 +727,8 @@
     }
 
     if (IsValidationEnabled(var->attributes, ast::DisabledValidation::kIgnoreAddressSpace) &&
-        (var->declared_address_space == type::AddressSpace::kIn ||
-         var->declared_address_space == type::AddressSpace::kOut)) {
+        (v->AddressSpace() == type::AddressSpace::kIn ||
+         v->AddressSpace() == type::AddressSpace::kOut)) {
         AddError("invalid use of input/output address space", var->source);
         return false;
     }
diff --git a/src/tint/transform/binding_remapper.cc b/src/tint/transform/binding_remapper.cc
index b657980..192b2c9 100644
--- a/src/tint/transform/binding_remapper.cc
+++ b/src/tint/transform/binding_remapper.cc
@@ -122,11 +122,12 @@
             // Replace any access controls.
             auto ac_it = remappings->access_controls.find(from);
             if (ac_it != remappings->access_controls.end()) {
-                type::Access ac = ac_it->second;
-                if (ac == type::Access::kUndefined) {
-                    b.Diagnostics().add_error(
-                        diag::System::Transform,
-                        "invalid access mode (" + std::to_string(static_cast<uint32_t>(ac)) + ")");
+                type::Access access = ac_it->second;
+                if (access == type::Access::kUndefined) {
+                    b.Diagnostics().add_error(diag::System::Transform,
+                                              "invalid access mode (" +
+                                                  std::to_string(static_cast<uint32_t>(access)) +
+                                                  ")");
                     return Program(std::move(b));
                 }
                 auto* sem = src->Sem().Get(var);
@@ -139,9 +140,14 @@
                 }
                 auto* ty = sem->Type()->UnwrapRef();
                 auto inner_ty = CreateASTTypeFor(ctx, ty);
-                auto* new_var = b.Var(ctx.Clone(var->source), ctx.Clone(var->name->symbol),
-                                      inner_ty, var->declared_address_space, ac,
-                                      ctx.Clone(var->initializer), ctx.Clone(var->attributes));
+                auto* new_var =
+                    b.create<ast::Var>(ctx.Clone(var->source),                  // source
+                                       b.Ident(ctx.Clone(var->name->symbol)),   // name
+                                       inner_ty,                                // type
+                                       ctx.Clone(var->declared_address_space),  // address space
+                                       b.Expr(access),                          // access
+                                       ctx.Clone(var->initializer),             // initializer
+                                       ctx.Clone(var->attributes));             // attributes
                 ctx.Replace(var, new_var);
             }
 
diff --git a/src/tint/transform/canonicalize_entry_point_io_test.cc b/src/tint/transform/canonicalize_entry_point_io_test.cc
index 5af8a14..f0ce8c3 100644
--- a/src/tint/transform/canonicalize_entry_point_io_test.cc
+++ b/src/tint/transform/canonicalize_entry_point_io_test.cc
@@ -67,11 +67,11 @@
 )";
 
     auto* expect = R"(
-@location(1) @internal(disable_validation__ignore_address_space) var<in> loc1_1 : f32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__in> loc1_1 : f32;
 
-@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> loc2_1 : vec4<u32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> loc2_1 : vec4<u32>;
 
-@builtin(position) @internal(disable_validation__ignore_address_space) var<in> coord_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__in> coord_1 : vec4<f32>;
 
 fn frag_main_inner(loc1 : f32, loc2 : vec4<u32>, coord : vec4<f32>) {
   var col : f32 = (coord.x * loc1);
@@ -251,13 +251,13 @@
 )";
 
     auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> loc0_1 : f32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__in> loc0_1 : f32;
 
-@location(1) @internal(disable_validation__ignore_address_space) var<in> loc1_1 : f32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__in> loc1_1 : f32;
 
-@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> loc2_1 : vec4<u32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> loc2_1 : vec4<u32>;
 
-@builtin(position) @internal(disable_validation__ignore_address_space) var<in> coord_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__in> coord_1 : vec4<f32>;
 
 struct FragBuiltins {
   coord : vec4<f32>,
@@ -304,13 +304,13 @@
 )";
 
     auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> loc0_1 : f32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__in> loc0_1 : f32;
 
-@location(1) @internal(disable_validation__ignore_address_space) var<in> loc1_1 : f32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__in> loc1_1 : f32;
 
-@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> loc2_1 : vec4<u32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> loc2_1 : vec4<u32>;
 
-@builtin(position) @internal(disable_validation__ignore_address_space) var<in> coord_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__in> coord_1 : vec4<f32>;
 
 fn frag_main_inner(loc0 : f32, locations : FragLocations, builtins : FragBuiltins) {
   var col : f32 = ((builtins.coord.x * locations.loc1) + loc0);
@@ -567,7 +567,7 @@
 )";
 
     auto* expect = R"(
-@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var<out> value : f32;
+@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var<__out> value : f32;
 
 fn frag_main_inner() -> f32 {
   return 1.0;
@@ -674,11 +674,11 @@
 )";
 
     auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<out> color_1 : vec4<f32>;
+@location(0) @internal(disable_validation__ignore_address_space) var<__out> color_1 : vec4<f32>;
 
-@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var<out> depth_1 : f32;
+@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var<__out> depth_1 : f32;
 
-@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<out> mask_1 : array<u32, 1u>;
+@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__out> mask_1 : array<u32, 1u>;
 
 struct FragOutput {
   color : vec4<f32>,
@@ -729,11 +729,11 @@
 )";
 
     auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<out> color_1 : vec4<f32>;
+@location(0) @internal(disable_validation__ignore_address_space) var<__out> color_1 : vec4<f32>;
 
-@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var<out> depth_1 : f32;
+@builtin(frag_depth) @internal(disable_validation__ignore_address_space) var<__out> depth_1 : f32;
 
-@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<out> mask_1 : array<u32, 1u>;
+@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__out> mask_1 : array<u32, 1u>;
 
 fn frag_main_inner() -> FragOutput {
   var output : FragOutput;
@@ -1028,13 +1028,13 @@
 )";
 
     auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> value_1 : f32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__in> value_1 : f32;
 
-@location(1) @internal(disable_validation__ignore_address_space) var<in> mul_1 : f32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__in> mul_1 : f32;
 
-@location(0) @internal(disable_validation__ignore_address_space) var<in> value_2 : f32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__in> value_2 : f32;
 
-@location(1) @internal(disable_validation__ignore_address_space) var<in> mul_2 : f32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__in> mul_2 : f32;
 
 struct FragmentInput {
   value : f32,
@@ -1094,13 +1094,13 @@
 )";
 
     auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> value_1 : f32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__in> value_1 : f32;
 
-@location(1) @internal(disable_validation__ignore_address_space) var<in> mul_1 : f32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__in> mul_1 : f32;
 
-@location(0) @internal(disable_validation__ignore_address_space) var<in> value_2 : f32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__in> value_2 : f32;
 
-@location(1) @internal(disable_validation__ignore_address_space) var<in> mul_2 : f32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__in> mul_2 : f32;
 
 fn frag_main1_inner(inputs : FragmentInput) {
   var x : f32 = foo(inputs);
@@ -1952,39 +1952,39 @@
 
     auto* expect =
         R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> i_1 : i32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__in> i_1 : i32;
 
-@location(1) @internal(disable_validation__ignore_address_space) var<in> u_1 : u32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__in> u_1 : u32;
 
-@location(2) @internal(disable_validation__ignore_address_space) var<in> vi_1 : vec4<i32>;
+@location(2) @internal(disable_validation__ignore_address_space) var<__in> vi_1 : vec4<i32>;
 
-@location(3) @internal(disable_validation__ignore_address_space) var<in> vu_1 : vec4<u32>;
+@location(3) @internal(disable_validation__ignore_address_space) var<__in> vu_1 : vec4<u32>;
 
-@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<out> i_2 : i32;
+@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> i_2 : i32;
 
-@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<out> u_2 : u32;
+@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> u_2 : u32;
 
-@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<out> vi_2 : vec4<i32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> vi_2 : vec4<i32>;
 
-@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<out> vu_2 : vec4<u32>;
+@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> vu_2 : vec4<u32>;
 
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> pos_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> pos_1 : vec4<f32>;
 
-@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> i_3 : i32;
+@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> i_3 : i32;
 
-@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> u_3 : u32;
+@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> u_3 : u32;
 
-@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> vi_3 : vec4<i32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> vi_3 : vec4<i32>;
 
-@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> vu_3 : vec4<u32>;
+@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> vu_3 : vec4<u32>;
 
-@location(0) @internal(disable_validation__ignore_address_space) var<out> i_4 : i32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__out> i_4 : i32;
 
-@location(1) @internal(disable_validation__ignore_address_space) var<out> u_4 : u32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__out> u_4 : u32;
 
-@location(2) @internal(disable_validation__ignore_address_space) var<out> vi_4 : vec4<i32>;
+@location(2) @internal(disable_validation__ignore_address_space) var<__out> vi_4 : vec4<i32>;
 
-@location(3) @internal(disable_validation__ignore_address_space) var<out> vu_4 : vec4<u32>;
+@location(3) @internal(disable_validation__ignore_address_space) var<__out> vu_4 : vec4<u32>;
 
 struct VertexIn {
   i : i32,
@@ -2082,39 +2082,39 @@
 
     auto* expect =
         R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> i_1 : i32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__in> i_1 : i32;
 
-@location(1) @internal(disable_validation__ignore_address_space) var<in> u_1 : u32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__in> u_1 : u32;
 
-@location(2) @internal(disable_validation__ignore_address_space) var<in> vi_1 : vec4<i32>;
+@location(2) @internal(disable_validation__ignore_address_space) var<__in> vi_1 : vec4<i32>;
 
-@location(3) @internal(disable_validation__ignore_address_space) var<in> vu_1 : vec4<u32>;
+@location(3) @internal(disable_validation__ignore_address_space) var<__in> vu_1 : vec4<u32>;
 
-@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<out> i_2 : i32;
+@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> i_2 : i32;
 
-@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<out> u_2 : u32;
+@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> u_2 : u32;
 
-@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<out> vi_2 : vec4<i32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> vi_2 : vec4<i32>;
 
-@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<out> vu_2 : vec4<u32>;
+@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__out> vu_2 : vec4<u32>;
 
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> pos_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> pos_1 : vec4<f32>;
 
-@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> i_3 : i32;
+@location(0) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> i_3 : i32;
 
-@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> u_3 : u32;
+@location(1) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> u_3 : u32;
 
-@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> vi_3 : vec4<i32>;
+@location(2) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> vi_3 : vec4<i32>;
 
-@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> vu_3 : vec4<u32>;
+@location(3) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> vu_3 : vec4<u32>;
 
-@location(0) @internal(disable_validation__ignore_address_space) var<out> i_4 : i32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__out> i_4 : i32;
 
-@location(1) @internal(disable_validation__ignore_address_space) var<out> u_4 : u32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__out> u_4 : u32;
 
-@location(2) @internal(disable_validation__ignore_address_space) var<out> vi_4 : vec4<i32>;
+@location(2) @internal(disable_validation__ignore_address_space) var<__out> vi_4 : vec4<i32>;
 
-@location(3) @internal(disable_validation__ignore_address_space) var<out> vu_4 : vec4<u32>;
+@location(3) @internal(disable_validation__ignore_address_space) var<__out> vu_4 : vec4<u32>;
 
 fn vert_main_inner(in : VertexIn) -> VertexOut {
   return VertexOut(in.i, in.u, in.vi, in.vu, vec4<f32>());
@@ -3161,9 +3161,9 @@
 )";
 
     auto* expect = R"(
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> value : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> value : vec4<f32>;
 
-@builtin(point_size) @internal(disable_validation__ignore_address_space) var<out> vertex_point_size : f32;
+@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size : f32;
 
 fn vert_main_inner() -> vec4<f32> {
   return vec4<f32>();
@@ -3236,9 +3236,9 @@
 )";
 
     auto* expect = R"(
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> pos_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> pos_1 : vec4<f32>;
 
-@builtin(point_size) @internal(disable_validation__ignore_address_space) var<out> vertex_point_size : f32;
+@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size : f32;
 
 struct VertOut {
   pos : vec4<f32>,
@@ -3277,9 +3277,9 @@
 )";
 
     auto* expect = R"(
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> pos_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> pos_1 : vec4<f32>;
 
-@builtin(point_size) @internal(disable_validation__ignore_address_space) var<out> vertex_point_size : f32;
+@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size : f32;
 
 fn vert_main_inner() -> VertOut {
   return VertOut();
@@ -3424,15 +3424,15 @@
 )";
 
     auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> collide_2 : f32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__in> collide_2 : f32;
 
-@location(1) @internal(disable_validation__ignore_address_space) var<in> collide_3 : f32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__in> collide_3 : f32;
 
-@location(0) @internal(disable_validation__ignore_address_space) var<out> vertex_point_size_3 : f32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_3 : f32;
 
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> vertex_point_size_1_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_1_1 : vec4<f32>;
 
-@builtin(point_size) @internal(disable_validation__ignore_address_space) var<out> vertex_point_size_4 : f32;
+@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_4 : f32;
 
 var<private> vertex_point_size : f32;
 
@@ -3502,15 +3502,15 @@
 )";
 
     auto* expect = R"(
-@location(0) @internal(disable_validation__ignore_address_space) var<in> collide_2 : f32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__in> collide_2 : f32;
 
-@location(1) @internal(disable_validation__ignore_address_space) var<in> collide_3 : f32;
+@location(1) @internal(disable_validation__ignore_address_space) var<__in> collide_3 : f32;
 
-@location(0) @internal(disable_validation__ignore_address_space) var<out> vertex_point_size_3 : f32;
+@location(0) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_3 : f32;
 
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> vertex_point_size_1_1 : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_1_1 : vec4<f32>;
 
-@builtin(point_size) @internal(disable_validation__ignore_address_space) var<out> vertex_point_size_4 : f32;
+@builtin(point_size) @internal(disable_validation__ignore_address_space) var<__out> vertex_point_size_4 : f32;
 
 fn vert_main_inner(collide : VertIn1, collide_1 : VertIn2) -> VertOut {
   let x = (collide.collide + collide_1.collide);
@@ -3868,11 +3868,11 @@
 )";
 
     auto* expect = R"(
-@builtin(sample_index) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> sample_index_1 : u32;
+@builtin(sample_index) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> sample_index_1 : u32;
 
-@builtin(sample_mask) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<in> mask_in_1 : array<u32, 1u>;
+@builtin(sample_mask) @interpolate(flat) @internal(disable_validation__ignore_address_space) var<__in> mask_in_1 : array<u32, 1u>;
 
-@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<out> value : array<u32, 1u>;
+@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__out> value : array<u32, 1u>;
 
 fn main_inner(sample_index : u32, mask_in : u32) -> u32 {
   return mask_in;
@@ -3903,11 +3903,11 @@
 )";
 
     auto* expect = R"(
-@builtin(sample_index) @internal(disable_validation__ignore_address_space) var<in> gl_SampleID : i32;
+@builtin(sample_index) @internal(disable_validation__ignore_address_space) var<__in> gl_SampleID : i32;
 
-@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<in> gl_SampleMaskIn : array<i32, 1u>;
+@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__in> gl_SampleMaskIn : array<i32, 1u>;
 
-@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<out> gl_SampleMask : array<i32, 1u>;
+@builtin(sample_mask) @internal(disable_validation__ignore_address_space) var<__out> gl_SampleMask : array<i32, 1u>;
 
 fn fragment_main(sample_index : u32, mask_in : u32) -> u32 {
   return mask_in;
@@ -3938,11 +3938,11 @@
 )";
 
     auto* expect = R"(
-@builtin(vertex_index) @internal(disable_validation__ignore_address_space) var<in> gl_VertexID : i32;
+@builtin(vertex_index) @internal(disable_validation__ignore_address_space) var<__in> gl_VertexID : i32;
 
-@builtin(instance_index) @internal(disable_validation__ignore_address_space) var<in> gl_InstanceID : i32;
+@builtin(instance_index) @internal(disable_validation__ignore_address_space) var<__in> gl_InstanceID : i32;
 
-@builtin(position) @internal(disable_validation__ignore_address_space) var<out> gl_Position : vec4<f32>;
+@builtin(position) @internal(disable_validation__ignore_address_space) var<__out> gl_Position : vec4<f32>;
 
 fn vertex_main(vertexID : u32, instanceID : u32) -> vec4<f32> {
   return vec4<f32>((f32(vertexID) + f32(instanceID)));
diff --git a/src/tint/transform/clamp_frag_depth.cc b/src/tint/transform/clamp_frag_depth.cc
index 57d9a40..8d84ed3 100644
--- a/src/tint/transform/clamp_frag_depth.cc
+++ b/src/tint/transform/clamp_frag_depth.cc
@@ -88,10 +88,11 @@
     // Abort on any use of push constants in the module.
     for (auto* global : src->AST().GlobalVariables()) {
         if (auto* var = global->As<ast::Var>()) {
-            if (TINT_UNLIKELY(var->declared_address_space == type::AddressSpace::kPushConstant)) {
+            auto* v = src->Sem().Get(var);
+            if (TINT_UNLIKELY(v->AddressSpace() == type::AddressSpace::kPushConstant)) {
                 TINT_ICE(Transform, b.Diagnostics())
                     << "ClampFragDepth doesn't know how to handle module that already use push "
-                       "constants.";
+                       "constants";
                 return Program(std::move(b));
             }
         }
diff --git a/src/tint/transform/module_scope_var_to_entry_point_param.cc b/src/tint/transform/module_scope_var_to_entry_point_param.cc
index ff65fec..47356e9 100644
--- a/src/tint/transform/module_scope_var_to_entry_point_param.cc
+++ b/src/tint/transform/module_scope_var_to_entry_point_param.cc
@@ -117,7 +117,6 @@
                                      WorkgroupParameterMemberList& workgroup_parameter_members,
                                      bool& is_pointer,
                                      bool& is_wrapped) {
-        auto* var_ast = var->Declaration()->As<ast::Var>();
         auto* ty = var->Type()->UnwrapRef();
 
         // Helper to create an AST node for the store type of the variable.
@@ -160,7 +159,9 @@
                     is_wrapped = true;
                 }
 
-                param_type = ctx.dst->ty.pointer(param_type, sc, var_ast->declared_access);
+                param_type = sc == type::AddressSpace::kStorage
+                                 ? ctx.dst->ty.pointer(param_type, sc, var->Access())
+                                 : ctx.dst->ty.pointer(param_type, sc);
                 auto* param = ctx.dst->Param(new_var_symbol, param_type, attributes);
                 ctx.InsertFront(func->params, param);
                 is_pointer = true;
@@ -228,7 +229,6 @@
                                        const sem::Variable* var,
                                        Symbol new_var_symbol,
                                        bool& is_pointer) {
-        auto* var_ast = var->Declaration()->As<ast::Var>();
         auto* ty = var->Type()->UnwrapRef();
         auto param_type = CreateASTTypeFor(ctx, ty);
         auto sc = var->AddressSpace();
@@ -254,7 +254,9 @@
         // Use a pointer for non-handle types.
         utils::Vector<const ast::Attribute*, 2> attributes;
         if (!ty->is_handle()) {
-            param_type = ctx.dst->ty.pointer(param_type, sc, var_ast->declared_access);
+            param_type = sc == type::AddressSpace::kStorage
+                             ? ctx.dst->ty.pointer(param_type, sc, var->Access())
+                             : ctx.dst->ty.pointer(param_type, sc);
             is_pointer = true;
 
             // Disable validation of the parameter's address space and of arguments passed to it.
diff --git a/src/tint/transform/module_scope_var_to_entry_point_param_test.cc b/src/tint/transform/module_scope_var_to_entry_point_param_test.cc
index 999761d..0f4c5b9 100644
--- a/src/tint/transform/module_scope_var_to_entry_point_param_test.cc
+++ b/src/tint/transform/module_scope_var_to_entry_point_param_test.cc
@@ -435,7 +435,7 @@
 }
 
 @compute @workgroup_size(1)
-fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, S>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, S, read>) {
   _ = *(tint_symbol);
   _ = *(tint_symbol_1);
 }
@@ -465,7 +465,7 @@
 
     auto* expect = R"(
 @compute @workgroup_size(1)
-fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, S>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, S, read>) {
   _ = *(tint_symbol);
   _ = *(tint_symbol_1);
 }
@@ -497,7 +497,7 @@
 }
 
 @compute @workgroup_size(1)
-fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1, read>) {
   _ = (*(tint_symbol)).arr[0];
 }
 )";
@@ -524,7 +524,7 @@
 }
 
 @compute @workgroup_size(1)
-fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1, read>) {
   _ = (*(tint_symbol)).arr[0];
 }
 )";
@@ -554,12 +554,12 @@
   arr : array<f32>,
 }
 
-fn foo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<storage, array<f32>>) {
+fn foo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<storage, array<f32>, read>) {
   _ = (*(tint_symbol))[0];
 }
 
 @compute @workgroup_size(1)
-fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, tint_symbol_2>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, tint_symbol_2, read>) {
   foo(&((*(tint_symbol_1)).arr));
 }
 )";
@@ -589,11 +589,11 @@
 }
 
 @compute @workgroup_size(1)
-fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, tint_symbol_2>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_1 : ptr<storage, tint_symbol_2, read>) {
   foo(&((*(tint_symbol_1)).arr));
 }
 
-fn foo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<storage, array<f32>>) {
+fn foo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<storage, array<f32>, read>) {
   _ = (*(tint_symbol))[0];
 }
 )";
@@ -624,7 +624,7 @@
 alias myarray = array<f32>;
 
 @compute @workgroup_size(1)
-fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1, read>) {
   _ = (*(tint_symbol)).arr[0];
 }
 )";
@@ -652,7 +652,7 @@
 }
 
 @compute @workgroup_size(1)
-fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1, read>) {
   _ = (*(tint_symbol)).arr[0];
 }
 
@@ -689,7 +689,7 @@
 }
 
 @compute @workgroup_size(1)
-fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1, read>) {
   _ = (*(tint_symbol)).arr[0];
 }
 )";
@@ -723,7 +723,7 @@
 }
 
 @compute @workgroup_size(1)
-fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol : ptr<storage, tint_symbol_1, read>) {
   _ = (*(tint_symbol)).arr[0];
 }
 )";
@@ -773,12 +773,12 @@
 fn no_uses() {
 }
 
-fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<storage, S>) {
+fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<storage, S, read>) {
   _ = *(tint_symbol);
   _ = *(tint_symbol_1);
 }
 
-fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr<storage, S>) {
+fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr<storage, S, read>) {
   let b : f32 = 2.0;
   _ = *(tint_symbol_2);
   bar(a, b, tint_symbol_2, tint_symbol_3);
@@ -786,7 +786,7 @@
 }
 
 @compute @workgroup_size(1)
-fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_4 : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_5 : ptr<storage, S>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_4 : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_5 : ptr<storage, S, read>) {
   foo(1.0, tint_symbol_4, tint_symbol_5);
 }
 )";
@@ -830,11 +830,11 @@
 
     auto* expect = R"(
 @compute @workgroup_size(1)
-fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_4 : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_5 : ptr<storage, S>) {
+fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_4 : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_address_space) tint_symbol_5 : ptr<storage, S, read>) {
   foo(1.0, tint_symbol_4, tint_symbol_5);
 }
 
-fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr<storage, S>) {
+fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr<storage, S, read>) {
   let b : f32 = 2.0;
   _ = *(tint_symbol_2);
   bar(a, b, tint_symbol_2, tint_symbol_3);
@@ -844,7 +844,7 @@
 fn no_uses() {
 }
 
-fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<storage, S>) {
+fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<uniform, S>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<storage, S, read>) {
   _ = *(tint_symbol);
   _ = *(tint_symbol_1);
 }
diff --git a/src/tint/transform/renamer_test.cc b/src/tint/transform/renamer_test.cc
index c233267..2d13436 100644
--- a/src/tint/transform/renamer_test.cc
+++ b/src/tint/transform/renamer_test.cc
@@ -1927,7 +1927,9 @@
         out.push_back(ident);
     }
     for (auto* ident : type::kAddressSpaceStrings) {
-        out.push_back(ident);
+        if (!utils::HasPrefix(ident, "_")) {
+            out.push_back(ident);
+        }
     }
     for (auto* ident : type::kTexelFormatStrings) {
         out.push_back(ident);
@@ -1940,24 +1942,24 @@
 
 using RenamerBuiltinIdentifierTest = TransformTestWithParam<const char*>;
 
-TEST_P(RenamerBuiltinIdentifierTest, GlobalVarName) {
+TEST_P(RenamerBuiltinIdentifierTest, GlobalConstName) {
     auto expand = [&](const char* source) {
         return utils::ReplaceAll(source, "$name", GetParam());
     };
 
     auto src = expand(R"(
-var<private> $name = 42;
+const $name = 42;
 
 fn f() {
-  var v = $name;
+  const v = $name;
 }
 )");
 
     auto expect = expand(R"(
-var<private> tint_symbol = 42;
+const tint_symbol = 42;
 
 fn tint_symbol_1() {
-  var tint_symbol_2 = tint_symbol;
+  const tint_symbol_2 = tint_symbol;
 }
 )");
 
diff --git a/src/tint/transform/std140.cc b/src/tint/transform/std140.cc
index 4c487d0..a19d78f 100644
--- a/src/tint/transform/std140.cc
+++ b/src/tint/transform/std140.cc
@@ -350,8 +350,8 @@
     void ReplaceUniformVarTypes() {
         for (auto* global : src->AST().GlobalVariables()) {
             if (auto* var = global->As<ast::Var>()) {
-                if (var->declared_address_space == type::AddressSpace::kUniform) {
-                    auto* v = sem.Get(var);
+                auto* v = sem.Get(var);
+                if (v->AddressSpace() == type::AddressSpace::kUniform) {
                     if (auto std140_ty = Std140Type(v->Type()->UnwrapRef())) {
                         ctx.Replace(global->type.expr, b.Expr(std140_ty));
                         std140_uniforms.Add(v);
diff --git a/src/tint/transform/zero_init_workgroup_memory.cc b/src/tint/transform/zero_init_workgroup_memory.cc
index 49d96d9..015a25b 100644
--- a/src/tint/transform/zero_init_workgroup_memory.cc
+++ b/src/tint/transform/zero_init_workgroup_memory.cc
@@ -36,7 +36,8 @@
 bool ShouldRun(const Program* program) {
     for (auto* global : program->AST().GlobalVariables()) {
         if (auto* var = global->As<ast::Var>()) {
-            if (var->declared_address_space == type::AddressSpace::kWorkgroup) {
+            auto* v = program->Sem().Get(var);
+            if (v->AddressSpace() == type::AddressSpace::kWorkgroup) {
                 return true;
             }
         }
diff --git a/src/tint/type/address_space.cc b/src/tint/type/address_space.cc
index 68f2c71..c3cfcc6 100644
--- a/src/tint/type/address_space.cc
+++ b/src/tint/type/address_space.cc
@@ -28,6 +28,12 @@
 /// @param str the string to parse
 /// @returns the parsed enum, or AddressSpace::kUndefined if the string could not be parsed.
 AddressSpace ParseAddressSpace(std::string_view str) {
+    if (str == "__in") {
+        return AddressSpace::kIn;
+    }
+    if (str == "__out") {
+        return AddressSpace::kOut;
+    }
     if (str == "function") {
         return AddressSpace::kFunction;
     }
@@ -53,14 +59,14 @@
     switch (value) {
         case AddressSpace::kUndefined:
             return out << "undefined";
+        case AddressSpace::kIn:
+            return out << "__in";
+        case AddressSpace::kOut:
+            return out << "__out";
         case AddressSpace::kFunction:
             return out << "function";
         case AddressSpace::kHandle:
             return out << "handle";
-        case AddressSpace::kIn:
-            return out << "in";
-        case AddressSpace::kOut:
-            return out << "out";
         case AddressSpace::kPrivate:
             return out << "private";
         case AddressSpace::kPushConstant:
diff --git a/src/tint/type/address_space.h b/src/tint/type/address_space.h
index 5b3855c..f0b0224 100644
--- a/src/tint/type/address_space.h
+++ b/src/tint/type/address_space.h
@@ -30,10 +30,10 @@
 /// Address space of a given pointer.
 enum class AddressSpace {
     kUndefined,
+    kIn,
+    kOut,
     kFunction,
     kHandle,  // Tint-internal enum entry - not parsed
-    kIn,      // Tint-internal enum entry - not parsed
-    kOut,     // Tint-internal enum entry - not parsed
     kPrivate,
     kPushConstant,
     kStorage,
@@ -52,7 +52,7 @@
 AddressSpace ParseAddressSpace(std::string_view str);
 
 constexpr const char* kAddressSpaceStrings[] = {
-    "function", "private", "push_constant", "storage", "uniform", "workgroup",
+    "__in", "__out", "function", "private", "push_constant", "storage", "uniform", "workgroup",
 };
 
 /// @returns true if the AddressSpace is host-shareable
diff --git a/src/tint/type/address_space_bench.cc b/src/tint/type/address_space_bench.cc
index ecd3ccf..6f33497 100644
--- a/src/tint/type/address_space_bench.cc
+++ b/src/tint/type/address_space_bench.cc
@@ -31,15 +31,62 @@
 
 void AddressSpaceParser(::benchmark::State& state) {
     const char* kStrings[] = {
-        "fccnctin",       "ucti3",         "functVon",      "function",      "1unction",
-        "unJtqqon",       "llun77tion",    "ppqqivtHH",     "prcv",          "bivaGe",
-        "private",        "priviive",      "8WWivate",      "pxxvate",       "pXh_cggnstant",
-        "pX_Vonstanu",    "push_consta3t", "push_constant", "push_constanE", "push_TTPnstant",
-        "puxxdh_constan", "s44orage",      "stSSraVVe",     "RtoR22e",       "storage",
-        "sFra9e",         "stoage",        "VOORRHge",      "unfoym",        "llnnrrf77rm",
-        "unif4r00",       "uniform",       "nfoom",         "zzform",        "uiiippo1",
-        "workgrouXX",     "wor55gro99nII", "wrrrkgroSSaHH", "workgroup",     "kkrHoup",
-        "jgkrouRR",       "wokroub",
+        "ccin",
+        "3",
+        "_Vin",
+        "__in",
+        "1_in",
+        "_qiJ",
+        "_lli77",
+        "__qHupp",
+        "vt",
+        "G_bt",
+        "__out",
+        "__viut",
+        "__8WWt",
+        "Mxxou",
+        "fuXggton",
+        "fuXtou",
+        "funct3on",
+        "function",
+        "funEtion",
+        "PPncTTion",
+        "xxuncddon",
+        "p44ivate",
+        "prSSvaVVe",
+        "RriR22e",
+        "private",
+        "pFva9e",
+        "priate",
+        "VOORRHte",
+        "push_constyn",
+        "punnh_crr77stallt",
+        "pu4h_cons00ant",
+        "push_constant",
+        "puoo_costan",
+        "ushzzcnstant",
+        "push_coii11apt",
+        "storaXXe",
+        "9II5tnnrage",
+        "stoaSSrHHYe",
+        "storage",
+        "stkke",
+        "jtogRa",
+        "sbrag",
+        "unifojm",
+        "niform",
+        "qform",
+        "uniform",
+        "uniNNrm",
+        "nifrvv",
+        "QQiform",
+        "workrorf",
+        "workjroup",
+        "wNNorkrou2",
+        "workgroup",
+        "workgrop",
+        "rrorkgroup",
+        "workgroGp",
     };
     for (auto _ : state) {
         for (auto* str : kStrings) {
diff --git a/src/tint/type/address_space_test.cc b/src/tint/type/address_space_test.cc
index 8b2af5c..7909f0e 100644
--- a/src/tint/type/address_space_test.cc
+++ b/src/tint/type/address_space_test.cc
@@ -42,6 +42,8 @@
 }
 
 static constexpr Case kValidCases[] = {
+    {"__in", AddressSpace::kIn},
+    {"__out", AddressSpace::kOut},
     {"function", AddressSpace::kFunction},
     {"private", AddressSpace::kPrivate},
     {"push_constant", AddressSpace::kPushConstant},
@@ -51,15 +53,18 @@
 };
 
 static constexpr Case kInvalidCases[] = {
-    {"fccnctin", AddressSpace::kUndefined},        {"ucti3", AddressSpace::kUndefined},
-    {"functVon", AddressSpace::kUndefined},        {"priv1te", AddressSpace::kUndefined},
-    {"pqiJate", AddressSpace::kUndefined},         {"privat7ll", AddressSpace::kUndefined},
-    {"pqqsh_pponstHnt", AddressSpace::kUndefined}, {"pus_cnstat", AddressSpace::kUndefined},
-    {"bus_Gonstant", AddressSpace::kUndefined},    {"storiive", AddressSpace::kUndefined},
-    {"8WWorage", AddressSpace::kUndefined},        {"sxxrage", AddressSpace::kUndefined},
-    {"uXforgg", AddressSpace::kUndefined},         {"nfoXm", AddressSpace::kUndefined},
-    {"unif3rm", AddressSpace::kUndefined},         {"workgroEp", AddressSpace::kUndefined},
-    {"woTTPkroup", AddressSpace::kUndefined},      {"ddorkroxxp", AddressSpace::kUndefined},
+    {"ccin", AddressSpace::kUndefined},          {"3", AddressSpace::kUndefined},
+    {"_Vin", AddressSpace::kUndefined},          {"__ou1", AddressSpace::kUndefined},
+    {"qq_Jt", AddressSpace::kUndefined},         {"__oll7t", AddressSpace::kUndefined},
+    {"qquntppHon", AddressSpace::kUndefined},    {"cnciv", AddressSpace::kUndefined},
+    {"funGion", AddressSpace::kUndefined},       {"priviive", AddressSpace::kUndefined},
+    {"8WWivate", AddressSpace::kUndefined},      {"pxxvate", AddressSpace::kUndefined},
+    {"pXh_cggnstant", AddressSpace::kUndefined}, {"pX_Vonstanu", AddressSpace::kUndefined},
+    {"push_consta3t", AddressSpace::kUndefined}, {"Etorage", AddressSpace::kUndefined},
+    {"sPTTrage", AddressSpace::kUndefined},      {"storadxx", AddressSpace::kUndefined},
+    {"u44iform", AddressSpace::kUndefined},      {"unSSfoVVm", AddressSpace::kUndefined},
+    {"RniR22m", AddressSpace::kUndefined},       {"w9rFroup", AddressSpace::kUndefined},
+    {"workgoup", AddressSpace::kUndefined},      {"woVROOrHup", AddressSpace::kUndefined},
 };
 
 using AddressSpaceParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/writer/wgsl/generator_impl.cc b/src/tint/writer/wgsl/generator_impl.cc
index ae34abb..c92e51a 100644
--- a/src/tint/writer/wgsl/generator_impl.cc
+++ b/src/tint/writer/wgsl/generator_impl.cc
@@ -34,8 +34,6 @@
 #include "src/tint/ast/workgroup_attribute.h"
 #include "src/tint/sem/struct.h"
 #include "src/tint/sem/switch_statement.h"
-#include "src/tint/type/access.h"
-#include "src/tint/type/texture_dimension.h"
 #include "src/tint/utils/math.h"
 #include "src/tint/utils/scoped_assignment.h"
 #include "src/tint/writer/float_to_string.h"
@@ -376,24 +374,6 @@
     return true;
 }
 
-bool GeneratorImpl::EmitAccess(std::ostream& out, const type::Access access) {
-    switch (access) {
-        case type::Access::kRead:
-            out << "read";
-            return true;
-        case type::Access::kWrite:
-            out << "write";
-            return true;
-        case type::Access::kReadWrite:
-            out << "read_write";
-            return true;
-        default:
-            break;
-    }
-    diagnostics_.add_error(diag::System::Writer, "unknown access");
-    return false;
-}
-
 bool GeneratorImpl::EmitStructType(const ast::Struct* str) {
     if (str->attributes.Length()) {
         if (!EmitAttributes(line(), str->attributes)) {
@@ -473,17 +453,18 @@
         v,  //
         [&](const ast::Var* var) {
             out << "var";
-            auto address_space = var->declared_address_space;
-            auto ac = var->declared_access;
-            if (address_space != type::AddressSpace::kUndefined || ac != type::Access::kUndefined) {
-                out << "<" << address_space;
-                if (ac != type::Access::kUndefined) {
+            if (var->declared_address_space || var->declared_access) {
+                out << "<";
+                TINT_DEFER(out << ">");
+                if (!EmitExpression(out, var->declared_address_space)) {
+                    return false;
+                }
+                if (var->declared_access) {
                     out << ", ";
-                    if (!EmitAccess(out, ac)) {
+                    if (!EmitExpression(out, var->declared_access)) {
                         return false;
                     }
                 }
-                out << ">";
             }
             return true;
         },
diff --git a/src/tint/writer/wgsl/generator_impl.h b/src/tint/writer/wgsl/generator_impl.h
index 56c0dfb..746ce0f 100644
--- a/src/tint/writer/wgsl/generator_impl.h
+++ b/src/tint/writer/wgsl/generator_impl.h
@@ -35,7 +35,6 @@
 #include "src/tint/ast/unary_op_expression.h"
 #include "src/tint/program.h"
 #include "src/tint/sem/struct.h"
-#include "src/tint/type/storage_texture.h"
 #include "src/tint/writer/text_generator.h"
 
 namespace tint::writer::wgsl {
@@ -209,11 +208,6 @@
     /// @param fmt the format to generate
     /// @returns true if the format is emitted
     bool EmitImageFormat(std::ostream& out, const type::TexelFormat fmt);
-    /// Handles emitting an access control
-    /// @param out the output stream
-    /// @param access the access to generate
-    /// @returns true if the access is emitted
-    bool EmitAccess(std::ostream& out, const type::Access access);
     /// Handles a unary op expression
     /// @param out the output stream
     /// @param expr the expression to emit