diff --git a/src/tint/cmd/common/generate_external_texture_bindings.cc b/src/tint/cmd/common/generate_external_texture_bindings.cc
index 7002faa..fdb7704 100644
--- a/src/tint/cmd/common/generate_external_texture_bindings.cc
+++ b/src/tint/cmd/common/generate_external_texture_bindings.cc
@@ -48,7 +48,7 @@
     std::vector<tint::BindingPoint> ext_tex_bps;
     for (auto* var : program.AST().GlobalVariables()) {
         if (auto* sem_var = program.Sem().Get(var)->As<sem::GlobalVariable>()) {
-            if (auto bp = sem_var->BindingPoint()) {
+            if (auto bp = sem_var->Attributes().binding_point) {
                 auto& n = group_to_next_binding_number[bp->group];
                 n = std::max(n, bp->binding + 1);
 
diff --git a/src/tint/fuzzers/tint_common_fuzzer.cc b/src/tint/fuzzers/tint_common_fuzzer.cc
index 8f36a89..e6ae17b 100644
--- a/src/tint/fuzzers/tint_common_fuzzer.cc
+++ b/src/tint/fuzzers/tint_common_fuzzer.cc
@@ -284,7 +284,7 @@
         std::vector<BindingPoint> ext_tex_bps;
         for (auto* var : program.AST().GlobalVariables()) {
             if (auto* sem_var = program.Sem().Get(var)->As<sem::GlobalVariable>()) {
-                if (auto bp = sem_var->BindingPoint()) {
+                if (auto bp = sem_var->Attributes().binding_point) {
                     auto& n = group_to_next_binding_number[bp->group];
                     n = std::max(n, bp->binding + 1);
 
diff --git a/src/tint/lang/core/ir/transform/shader_io.cc b/src/tint/lang/core/ir/transform/shader_io.cc
index 0fd0ac8..368dcf0 100644
--- a/src/tint/lang/core/ir/transform/shader_io.cc
+++ b/src/tint/lang/core/ir/transform/shader_io.cc
@@ -119,7 +119,13 @@
         if (func->Stage() == Function::PipelineStage::kVertex && backend->NeedsVertexPointSize()) {
             vertex_point_size_index =
                 backend->AddOutput(ir.symbols.New("vertex_point_size"), ty.f32(),
-                                   {{}, {}, {BuiltinValue::kPointSize}, {}, false});
+                                   core::type::StructMemberAttributes{
+                                       /* location */ {},
+                                       /* index */ {},
+                                       /* builtin */ core::BuiltinValue::kPointSize,
+                                       /* interpolation */ {},
+                                       /* invariant */ false,
+                                   });
         }
 
         auto new_params = backend->FinalizeInputs();
diff --git a/src/tint/lang/core/ir/transform/zero_init_workgroup_memory_test.cc b/src/tint/lang/core/ir/transform/zero_init_workgroup_memory_test.cc
index 0cf6dc4..e16f35e 100644
--- a/src/tint/lang/core/ir/transform/zero_init_workgroup_memory_test.cc
+++ b/src/tint/lang/core/ir/transform/zero_init_workgroup_memory_test.cc
@@ -1486,20 +1486,31 @@
 TEST_F(IR_ZeroInitWorkgroupMemoryTest, ExistingLocalInvocationIndexInStruct) {
     auto* var = MakeVar("wgvar", ty.bool_());
 
-    auto* structure =
-        ty.Struct(mod.symbols.New("MyStruct"),
-                  {
-                      {
-                          mod.symbols.New("global_id"),
-                          ty.vec3<u32>(),
-                          {{}, {}, core::BuiltinValue::kGlobalInvocationId, {}, false},
-                      },
-                      {
-                          mod.symbols.New("index"),
-                          ty.u32(),
-                          {{}, {}, core::BuiltinValue::kLocalInvocationIndex, {}, false},
-                      },
-                  });
+    auto* structure = ty.Struct(mod.symbols.New("MyStruct"),
+                                {
+                                    {
+                                        mod.symbols.New("global_id"),
+                                        ty.vec3<u32>(),
+                                        core::type::StructMemberAttributes{
+                                            /* location */ {},
+                                            /* index */ {},
+                                            /* builtin */ core::BuiltinValue::kGlobalInvocationId,
+                                            /* interpolation */ {},
+                                            /* invariant */ false,
+                                        },
+                                    },
+                                    {
+                                        mod.symbols.New("index"),
+                                        ty.u32(),
+                                        core::type::StructMemberAttributes{
+                                            /* location */ {},
+                                            /* index */ {},
+                                            /* builtin */ core::BuiltinValue::kLocalInvocationIndex,
+                                            /* interpolation */ {},
+                                            /* invariant */ false,
+                                        },
+                                    },
+                                });
     auto* func = MakeEntryPoint("main", 1, 1, 1);
     func->SetParams({b.FunctionParam("params", structure)});
     b.Append(func->Block(), [&] {  //
diff --git a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
index b0c7219..738844e 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
@@ -1929,7 +1929,7 @@
         TINT_ICE() << "storage variable must be of struct type";
         return;
     }
-    auto bp = *sem->As<sem::GlobalVariable>()->BindingPoint();
+    auto bp = *sem->As<sem::GlobalVariable>()->Attributes().binding_point;
     {
         auto out = Line();
         out << "layout(binding = " << bp.binding << ", std140";
@@ -1948,7 +1948,7 @@
         TINT_ICE() << "storage variable must be of struct type";
         return;
     }
-    auto bp = *sem->As<sem::GlobalVariable>()->BindingPoint();
+    auto bp = *sem->As<sem::GlobalVariable>()->Attributes().binding_point;
     Line() << "layout(binding = " << bp.binding << ", std430) buffer "
            << UniqueIdentifier(StructName(str) + "_ssbo") << " {";
     EmitStructMembers(current_buffer_, str);
@@ -2080,7 +2080,7 @@
     }
 
     auto out = Line();
-    EmitAttributes(out, var, decl->attributes);
+    EmitAttributes(out, var);
     EmitInterpolationQualifiers(out, decl->attributes);
 
     auto name = decl->name->symbol.Name();
@@ -2130,23 +2130,17 @@
     }
 }
 
-void ASTPrinter::EmitAttributes(StringStream& out,
-                                const sem::GlobalVariable* var,
-                                VectorRef<const ast::Attribute*> attributes) {
-    if (attributes.IsEmpty()) {
-        return;
-    }
+void ASTPrinter::EmitAttributes(StringStream& out, const sem::GlobalVariable* var) {
+    auto& attrs = var->Attributes();
 
     bool first = true;
-    for (auto* attr : attributes) {
-        if (attr->As<ast::LocationAttribute>()) {
-            out << (first ? "layout(" : ", ");
-            out << "location = " << std::to_string(var->Location().value());
-            first = false;
-        }
-        if (attr->As<ast::IndexAttribute>()) {
-            out << ", index = " << std::to_string(var->Index().value());
-        }
+    if (attrs.location.has_value()) {
+        out << (first ? "layout(" : ", ");
+        out << "location = " << std::to_string(attrs.location.value());
+        first = false;
+    }
+    if (attrs.index.has_value()) {
+        out << ", index = " << std::to_string(attrs.index.value());
     }
     if (!first) {
         out << ") ";
diff --git a/src/tint/lang/glsl/writer/ast_printer/ast_printer.h b/src/tint/lang/glsl/writer/ast_printer/ast_printer.h
index 519bdf8..f5cc633 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.h
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.h
@@ -317,10 +317,7 @@
     /// Handles emitting attributes
     /// @param out the output of the expression stream
     /// @param var the global variable semantics
-    /// @param attrs the attributes
-    void EmitAttributes(StringStream& out,
-                        const sem::GlobalVariable* var,
-                        VectorRef<const ast::Attribute*> attrs);
+    void EmitAttributes(StringStream& out, const sem::GlobalVariable* var);
     /// Handles emitting the entry point function
     /// @param func the entry point
     void EmitEntryPointFunction(const ast::Function* func);
diff --git a/src/tint/lang/glsl/writer/ast_raise/combine_samplers.cc b/src/tint/lang/glsl/writer/ast_raise/combine_samplers.cc
index b41e960..b056310 100644
--- a/src/tint/lang/glsl/writer/ast_raise/combine_samplers.cc
+++ b/src/tint/lang/glsl/writer/ast_raise/combine_samplers.cc
@@ -118,10 +118,10 @@
                                               std::string name) {
         SamplerTexturePair bp_pair;
         bp_pair.texture_binding_point =
-            texture_var ? *texture_var->As<sem::GlobalVariable>()->BindingPoint()
+            texture_var ? *texture_var->As<sem::GlobalVariable>()->Attributes().binding_point
                         : binding_info->placeholder_binding_point;
         bp_pair.sampler_binding_point =
-            sampler_var ? *sampler_var->As<sem::GlobalVariable>()->BindingPoint()
+            sampler_var ? *sampler_var->As<sem::GlobalVariable>()->Attributes().binding_point
                         : binding_info->placeholder_binding_point;
         auto it = binding_info->binding_map.find(bp_pair);
         if (it != binding_info->binding_map.end()) {
@@ -232,7 +232,7 @@
             if (tint::IsAnyOf<core::type::Texture, core::type::Sampler>(type) &&
                 !type->Is<core::type::StorageTexture>()) {
                 ctx.Remove(ctx.src->AST().GlobalDeclarations(), global);
-            } else if (auto binding_point = global_sem->BindingPoint()) {
+            } else if (auto binding_point = global_sem->Attributes().binding_point) {
                 if (binding_point->group == 0 && binding_point->binding == 0) {
                     auto* attribute =
                         ctx.dst->Disable(ast::DisabledValidation::kBindingPointCollision);
diff --git a/src/tint/lang/glsl/writer/ast_raise/texture_builtins_from_uniform.cc b/src/tint/lang/glsl/writer/ast_raise/texture_builtins_from_uniform.cc
index 67eb338..bc40549 100644
--- a/src/tint/lang/glsl/writer/ast_raise/texture_builtins_from_uniform.cc
+++ b/src/tint/lang/glsl/writer/ast_raise/texture_builtins_from_uniform.cc
@@ -366,7 +366,7 @@
                 auto* global_sem = sem.Get<sem::GlobalVariable>(var);
 
                 // The original binding point
-                BindingPoint binding_point = *global_sem->BindingPoint();
+                BindingPoint binding_point = *global_sem->Attributes().binding_point;
 
                 if (binding_point == cfg->ubo_binding) {
                     // This ubo_binding struct already exists.
@@ -423,7 +423,7 @@
     /// @returns binding of the global variable.
     BindingPoint GetAndRecordGlobalBinding(const sem::GlobalVariable* global,
                                            TextureBuiltinsFromUniformOptions::Field field) {
-        auto binding = global->BindingPoint().value();
+        auto binding = global->Attributes().binding_point.value();
         auto iter = bindpoint_to_data.find(binding);
         if (iter == bindpoint_to_data.end()) {
             // First visit, recording the binding.
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
index b23c19f..6ee9b5a 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -3331,7 +3331,7 @@
 }
 
 bool ASTPrinter::EmitUniformVariable(const ast::Var* var, const sem::Variable* sem) {
-    auto binding_point = *sem->As<sem::GlobalVariable>()->BindingPoint();
+    auto binding_point = *sem->As<sem::GlobalVariable>()->Attributes().binding_point;
     auto* type = sem->Type()->UnwrapRef();
     auto name = var->name->symbol.Name();
     Line() << "cbuffer cbuffer_" << name << RegisterAndSpace('b', binding_point) << " {";
@@ -3360,7 +3360,7 @@
 
     auto* global_sem = sem->As<sem::GlobalVariable>();
     out << RegisterAndSpace(sem->Access() == core::Access::kRead ? 't' : 'u',
-                            *global_sem->BindingPoint())
+                            *global_sem->Attributes().binding_point)
         << ";";
 
     return true;
@@ -3389,7 +3389,7 @@
     }
 
     if (register_space) {
-        auto bp = sem->As<sem::GlobalVariable>()->BindingPoint();
+        auto bp = sem->As<sem::GlobalVariable>()->Attributes().binding_point;
         out << " : register(" << register_space << bp->binding;
         // Omit the space if it's 0, as it's the default.
         // SM 5.0 doesn't support spaces, so we don't emit them if group is 0 for better
diff --git a/src/tint/lang/hlsl/writer/ast_raise/num_workgroups_from_uniform.cc b/src/tint/lang/hlsl/writer/ast_raise/num_workgroups_from_uniform.cc
index 31590d9..22a8ff3 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/num_workgroups_from_uniform.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/num_workgroups_from_uniform.cc
@@ -165,7 +165,7 @@
 
                 for (auto* global : src.AST().GlobalVariables()) {
                     auto* global_sem = src.Sem().Get<sem::GlobalVariable>(global);
-                    if (auto bp = global_sem->BindingPoint()) {
+                    if (auto bp = global_sem->Attributes().binding_point) {
                         if (bp->group >= group) {
                             group = bp->group + 1;
                         }
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index aed7504..6ef2020 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -1970,10 +1970,9 @@
             return kInvalidBindingIndex;
         }
         auto* param_sem = builder_.Sem().Get<sem::Parameter>(param);
-        auto bp = param_sem->BindingPoint();
+        auto bp = param_sem->Attributes().binding_point;
         if (TINT_UNLIKELY(bp->group != 0)) {
-            TINT_ICE() << "encountered non-zero resource group index (use "
-                          "BindingRemapper to fix)";
+            TINT_ICE() << "encountered non-zero resource group index (use BindingRemapper to fix)";
             return kInvalidBindingIndex;
         }
         return bp->binding;
diff --git a/src/tint/lang/msl/writer/writer_bench.cc b/src/tint/lang/msl/writer/writer_bench.cc
index 2705e93..0af1f7b 100644
--- a/src/tint/lang/msl/writer/writer_bench.cc
+++ b/src/tint/lang/msl/writer/writer_bench.cc
@@ -65,7 +65,7 @@
     uint32_t next_binding_point = 0;
     for (auto* var : program.AST().GlobalVariables()) {
         if (auto* var_sem = program.Sem().Get(var)->As<sem::GlobalVariable>()) {
-            if (auto bp = var_sem->BindingPoint()) {
+            if (auto bp = var_sem->Attributes().binding_point) {
                 gen_options.binding_remapper_options.binding_points[*bp] = BindingPoint{
                     0,                     // group
                     next_binding_point++,  // binding
diff --git a/src/tint/lang/spirv/writer/ast_printer/builder.cc b/src/tint/lang/spirv/writer/ast_printer/builder.cc
index 734c380..a75d703 100644
--- a/src/tint/lang/spirv/writer/ast_printer/builder.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/builder.cc
@@ -806,13 +806,13 @@
             [&](const ast::LocationAttribute*) {
                 module_.PushAnnot(spv::Op::OpDecorate,
                                   {Operand(var_id), U32Operand(SpvDecorationLocation),
-                                   Operand(sem->Location().value())});
+                                   Operand(sem->Attributes().location.value())});
                 return true;
             },
             [&](const ast::IndexAttribute*) {
                 module_.PushAnnot(spv::Op::OpDecorate,
                                   {Operand(var_id), U32Operand(SpvDecorationIndex),
-                                   Operand(sem->Index().value())});
+                                   Operand(sem->Attributes().index.value())});
                 return true;
             },
             [&](const ast::InterpolateAttribute* interpolate) {
@@ -837,14 +837,14 @@
                 return true;
             },
             [&](const ast::BindingAttribute*) {
-                auto bp = sem->BindingPoint();
+                auto bp = sem->Attributes().binding_point;
                 module_.PushAnnot(
                     spv::Op::OpDecorate,
                     {Operand(var_id), U32Operand(SpvDecorationBinding), Operand(bp->binding)});
                 return true;
             },
             [&](const ast::GroupAttribute*) {
-                auto bp = sem->BindingPoint();
+                auto bp = sem->Attributes().binding_point;
                 module_.PushAnnot(
                     spv::Op::OpDecorate,
                     {Operand(var_id), U32Operand(SpvDecorationDescriptorSet), Operand(bp->group)});
diff --git a/src/tint/lang/spirv/writer/function_test.cc b/src/tint/lang/spirv/writer/function_test.cc
index 890ddff..fc66154 100644
--- a/src/tint/lang/spirv/writer/function_test.cc
+++ b/src/tint/lang/spirv/writer/function_test.cc
@@ -347,11 +347,31 @@
 }
 
 TEST_F(SpirvWriterTest, Function_ShaderIO_DualSourceBlend) {
-    auto* outputs = ty.Struct(mod.symbols.New("Outputs"),
-                              {
-                                  {mod.symbols.Register("a"), ty.f32(), {0u, 0u, {}, {}, false}},
-                                  {mod.symbols.Register("b"), ty.f32(), {0u, 1u, {}, {}, false}},
-                              });
+    auto* outputs =
+        ty.Struct(mod.symbols.New("Outputs"), {
+                                                  {
+                                                      mod.symbols.Register("a"),
+                                                      ty.f32(),
+                                                      core::type::StructMemberAttributes{
+                                                          /* location */ 0u,
+                                                          /* index */ 0u,
+                                                          /* builtin */ std::nullopt,
+                                                          /* interpolation */ std::nullopt,
+                                                          /* invariant */ false,
+                                                      },
+                                                  },
+                                                  {
+                                                      mod.symbols.Register("b"),
+                                                      ty.f32(),
+                                                      core::type::StructMemberAttributes{
+                                                          /* location */ 0u,
+                                                          /* index */ 1u,
+                                                          /* builtin */ std::nullopt,
+                                                          /* interpolation */ std::nullopt,
+                                                          /* invariant */ false,
+                                                      },
+                                                  },
+                                              });
 
     auto* func = b.Function("main", outputs, core::ir::Function::PipelineStage::kFragment);
     b.Append(func->Block(), [&] {  //
diff --git a/src/tint/lang/spirv/writer/helpers/generate_bindings.cc b/src/tint/lang/spirv/writer/helpers/generate_bindings.cc
index 21ef88b..076a594 100644
--- a/src/tint/lang/spirv/writer/helpers/generate_bindings.cc
+++ b/src/tint/lang/spirv/writer/helpers/generate_bindings.cc
@@ -55,7 +55,7 @@
     Vector<tint::BindingPoint, 4> ext_tex_bps;
     for (auto* var : program.AST().GlobalVariables()) {
         if (auto* sem_var = program.Sem().Get(var)->As<sem::GlobalVariable>()) {
-            if (auto bp = sem_var->BindingPoint()) {
+            if (auto bp = sem_var->Attributes().binding_point) {
                 if (auto val = group_to_next_binding_number.Find(bp->group)) {
                     *val = std::max(*val, bp->binding + 1);
                 } else {
diff --git a/src/tint/lang/spirv/writer/raise/shader_io_test.cc b/src/tint/lang/spirv/writer/raise/shader_io_test.cc
index 4dcd4a3..edbde41 100644
--- a/src/tint/lang/spirv/writer/raise/shader_io_test.cc
+++ b/src/tint/lang/spirv/writer/raise/shader_io_test.cc
@@ -151,27 +151,50 @@
                                  {
                                      mod.symbols.New("front_facing"),
                                      ty.bool_(),
-                                     {{}, {}, core::BuiltinValue::kFrontFacing, {}, false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ std::nullopt,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ core::BuiltinValue::kFrontFacing,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ false,
+                                     },
                                  },
                                  {
                                      mod.symbols.New("position"),
                                      ty.vec4<f32>(),
-                                     {{}, {}, core::BuiltinValue::kPosition, {}, true},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ std::nullopt,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ core::BuiltinValue::kPosition,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ true,
+                                     },
                                  },
                                  {
                                      mod.symbols.New("color1"),
                                      ty.f32(),
-                                     {0u, {}, {}, {}, false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ 0u,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ std::nullopt,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ false,
+                                     },
                                  },
                                  {
                                      mod.symbols.New("color2"),
                                      ty.f32(),
-                                     {1u,
-                                      {},
-                                      {},
-                                      core::Interpolation{core::InterpolationType::kLinear,
-                                                          core::InterpolationSampling::kSample},
-                                      false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ 1u,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ std::nullopt,
+                                         /* interpolation */
+                                         core::Interpolation{
+                                             core::InterpolationType::kLinear,
+                                             core::InterpolationSampling::kSample,
+                                         },
+                                         /* invariant */ false,
+                                     },
                                  },
                              });
 
@@ -276,12 +299,24 @@
                                  {
                                      mod.symbols.New("position"),
                                      ty.vec4<f32>(),
-                                     {{}, {}, core::BuiltinValue::kPosition, {}, true},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ std::nullopt,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ core::BuiltinValue::kPosition,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ true,
+                                     },
                                  },
                                  {
                                      mod.symbols.New("color1"),
                                      ty.f32(),
-                                     {0u, {}, {}, {}, false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ 0u,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ std::nullopt,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ false,
+                                     },
                                  },
                              });
 
@@ -476,22 +511,39 @@
                                  {
                                      mod.symbols.New("position"),
                                      ty.vec4<f32>(),
-                                     {{}, {}, core::BuiltinValue::kPosition, {}, true},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ std::nullopt,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ core::BuiltinValue::kPosition,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ true,
+                                     },
                                  },
                                  {
                                      mod.symbols.New("color1"),
                                      ty.f32(),
-                                     {0u, {}, {}, {}, false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ 0u,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ std::nullopt,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ false,
+                                     },
                                  },
                                  {
                                      mod.symbols.New("color2"),
                                      ty.f32(),
-                                     {1u,
-                                      {},
-                                      {},
-                                      core::Interpolation{core::InterpolationType::kLinear,
-                                                          core::InterpolationSampling::kSample},
-                                      false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ 1u,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ std::nullopt,
+                                         /* interpolation */
+                                         core::Interpolation{
+                                             core::InterpolationType::kLinear,
+                                             core::InterpolationSampling::kSample,
+                                         },
+                                         /* invariant */ false,
+                                     },
                                  },
                              });
 
@@ -561,18 +613,31 @@
 }
 
 TEST_F(SpirvWriter_ShaderIOTest, ReturnValue_DualSourceBlending) {
-    auto* str_ty = ty.Struct(mod.symbols.New("Output"), {
-                                                            {
-                                                                mod.symbols.New("color1"),
-                                                                ty.f32(),
-                                                                {0u, 0u, {}, {}, false},
-                                                            },
-                                                            {
-                                                                mod.symbols.New("color2"),
-                                                                ty.f32(),
-                                                                {0u, 1u, {}, {}, false},
-                                                            },
-                                                        });
+    auto* str_ty =
+        ty.Struct(mod.symbols.New("Output"), {
+                                                 {
+                                                     mod.symbols.New("color1"),
+                                                     ty.f32(),
+                                                     core::type::StructMemberAttributes{
+                                                         /* location */ 0u,
+                                                         /* index */ 0u,
+                                                         /* builtin */ std::nullopt,
+                                                         /* interpolation */ std::nullopt,
+                                                         /* invariant */ false,
+                                                     },
+                                                 },
+                                                 {
+                                                     mod.symbols.New("color2"),
+                                                     ty.f32(),
+                                                     core::type::StructMemberAttributes{
+                                                         /* location */ 0u,
+                                                         /* index */ 1u,
+                                                         /* builtin */ std::nullopt,
+                                                         /* interpolation */ std::nullopt,
+                                                         /* invariant */ false,
+                                                     },
+                                                 },
+                                             });
 
     auto* ep = b.Function("foo", str_ty);
     ep->SetStage(core::ir::Function::PipelineStage::kFragment);
@@ -639,12 +704,24 @@
                                  {
                                      mod.symbols.New("position"),
                                      vec4f,
-                                     {{}, {}, core::BuiltinValue::kPosition, {}, false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ std::nullopt,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ core::BuiltinValue::kPosition,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ false,
+                                     },
                                  },
                                  {
                                      mod.symbols.New("color"),
                                      vec4f,
-                                     {0u, {}, {}, {}, false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ 0u,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ std::nullopt,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ false,
+                                     },
                                  },
                              });
 
@@ -766,12 +843,24 @@
                                  {
                                      mod.symbols.New("position"),
                                      vec4f,
-                                     {{}, {}, core::BuiltinValue::kPosition, {}, false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ std::nullopt,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ core::BuiltinValue::kPosition,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ false,
+                                     },
                                  },
                                  {
                                      mod.symbols.New("color"),
                                      vec4f,
-                                     {0u, {}, {}, {}, false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ 0u,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ std::nullopt,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ false,
+                                     },
                                  },
                              });
 
@@ -847,12 +936,24 @@
                                  {
                                      mod.symbols.New("color"),
                                      ty.f32(),
-                                     {0u, {}, {}, {}, false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ 0u,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ std::nullopt,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ false,
+                                     },
                                  },
                                  {
                                      mod.symbols.New("mask"),
                                      ty.u32(),
-                                     {{}, {}, core::BuiltinValue::kSampleMask, {}, false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ std::nullopt,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ core::BuiltinValue::kSampleMask,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ false,
+                                     },
                                  },
                              });
 
@@ -924,19 +1025,24 @@
 
 // Test that interpolation attributes are stripped from vertex inputs and fragment outputs.
 TEST_F(SpirvWriter_ShaderIOTest, InterpolationOnVertexInputOrFragmentOutput) {
-    auto* str_ty = ty.Struct(mod.symbols.New("MyStruct"),
-                             {
-                                 {
-                                     mod.symbols.New("color"),
-                                     ty.f32(),
-                                     {1u,
-                                      {},
-                                      {},
-                                      core::Interpolation{core::InterpolationType::kLinear,
-                                                          core::InterpolationSampling::kSample},
-                                      false},
-                                 },
-                             });
+    auto* str_ty =
+        ty.Struct(mod.symbols.New("MyStruct"), {
+                                                   {
+                                                       mod.symbols.New("color"),
+                                                       ty.f32(),
+                                                       core::type::StructMemberAttributes{
+                                                           /* location */ 1u,
+                                                           /* index */ std::nullopt,
+                                                           /* builtin */ std::nullopt,
+                                                           /* interpolation */
+                                                           core::Interpolation{
+                                                               core::InterpolationType::kLinear,
+                                                               core::InterpolationSampling::kSample,
+                                                           },
+                                                           /* invariant */ false,
+                                                       },
+                                                   },
+                                               });
 
     // Vertex shader.
     {
@@ -1071,12 +1177,24 @@
                                  {
                                      mod.symbols.New("color"),
                                      ty.f32(),
-                                     {0u, {}, {}, {}, false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ 0u,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ std::nullopt,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ false,
+                                     },
                                  },
                                  {
                                      mod.symbols.New("depth"),
                                      ty.f32(),
-                                     {{}, {}, core::BuiltinValue::kFragDepth, {}, false},
+                                     core::type::StructMemberAttributes{
+                                         /* location */ std::nullopt,
+                                         /* index */ std::nullopt,
+                                         /* builtin */ core::BuiltinValue::kFragDepth,
+                                         /* interpolation */ std::nullopt,
+                                         /* invariant */ false,
+                                     },
                                  },
                              });
 
diff --git a/src/tint/lang/wgsl/ast/transform/array_length_from_uniform.cc b/src/tint/lang/wgsl/ast/transform/array_length_from_uniform.cc
index 2f617c1..d98068d 100644
--- a/src/tint/lang/wgsl/ast/transform/array_length_from_uniform.cc
+++ b/src/tint/lang/wgsl/ast/transform/array_length_from_uniform.cc
@@ -100,7 +100,7 @@
 
         IterateArrayLengthOnStorageVar(
             [&](const CallExpression*, const sem::VariableUser*, const sem::GlobalVariable* var) {
-                if (auto binding = var->BindingPoint()) {
+                if (auto binding = var->Attributes().binding_point) {
                     auto idx_itr = cfg->bindpoint_to_size_index.find(*binding);
                     if (idx_itr == cfg->bindpoint_to_size_index.end()) {
                         return;
@@ -138,7 +138,7 @@
         IterateArrayLengthOnStorageVar([&](const CallExpression* call_expr,
                                            const sem::VariableUser* storage_buffer_sem,
                                            const sem::GlobalVariable* var) {
-            auto binding = var->BindingPoint();
+            auto binding = var->Attributes().binding_point;
             if (!binding) {
                 return;
             }
diff --git a/src/tint/lang/wgsl/ast/transform/binding_remapper.cc b/src/tint/lang/wgsl/ast/transform/binding_remapper.cc
index 3ff5e0a..b2ec92a 100644
--- a/src/tint/lang/wgsl/ast/transform/binding_remapper.cc
+++ b/src/tint/lang/wgsl/ast/transform/binding_remapper.cc
@@ -89,7 +89,7 @@
             auto* func = src.Sem().Get(func_ast);
             std::unordered_map<BindingPoint, int> binding_point_counts;
             for (auto* global : func->TransitivelyReferencedGlobals()) {
-                if (auto from = global->BindingPoint()) {
+                if (auto from = global->Attributes().binding_point) {
                     auto bp_it = remappings->binding_points.find(*from);
                     if (bp_it != remappings->binding_points.end()) {
                         // Remapped
@@ -113,7 +113,7 @@
             auto* global_sem = src.Sem().Get<sem::GlobalVariable>(var);
 
             // The original binding point
-            BindingPoint from = *global_sem->BindingPoint();
+            BindingPoint from = *global_sem->Attributes().binding_point;
 
             // The binding point after remapping
             BindingPoint bp = from;
diff --git a/src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io.cc b/src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io.cc
index 4643176..45c270c 100644
--- a/src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io.cc
+++ b/src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io.cc
@@ -428,7 +428,8 @@
         }
 
         auto name = param->Declaration()->name->symbol.Name();
-        auto* input_expr = AddInput(name, param->Type(), param->Location(), std::move(attributes));
+        auto* input_expr =
+            AddInput(name, param->Type(), param->Attributes().location, std::move(attributes));
         inner_call_parameters.Push(input_expr);
     }
 
diff --git a/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.cc b/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.cc
index c463dbc..0a5f85e 100644
--- a/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.cc
+++ b/src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.cc
@@ -131,7 +131,7 @@
             // The binding points for the newly introduced bindings must have been provided to this
             // transform. We fetch the new binding points by providing the original texture_external
             // binding points into the passed map.
-            BindingPoint bp = *sem_var->BindingPoint();
+            BindingPoint bp = *sem_var->Attributes().binding_point;
 
             BindingsMap::const_iterator it = new_binding_points->bindings_map.find(bp);
             if (it == new_binding_points->bindings_map.end()) {
diff --git a/src/tint/lang/wgsl/ast/transform/robustness.cc b/src/tint/lang/wgsl/ast/transform/robustness.cc
index c984b29..0bc1575 100644
--- a/src/tint/lang/wgsl/ast/transform/robustness.cc
+++ b/src/tint/lang/wgsl/ast/transform/robustness.cc
@@ -699,11 +699,11 @@
         if (globalVariable == nullptr) {
             return false;
         }
-        if (!globalVariable->BindingPoint().has_value()) {
+        auto binding_point = globalVariable->Attributes().binding_point;
+        if (!binding_point.has_value()) {
             return false;
         }
-        BindingPoint bindingPoint = *globalVariable->BindingPoint();
-        return cfg.bindings_ignored.find(bindingPoint) != cfg.bindings_ignored.cend();
+        return cfg.bindings_ignored.find(*binding_point) != cfg.bindings_ignored.cend();
     }
 
     /// @returns true if expr is an IndexAccessorExpression whose object is a runtime-sized array.
diff --git a/src/tint/lang/wgsl/ast/transform/single_entry_point.cc b/src/tint/lang/wgsl/ast/transform/single_entry_point.cc
index daae6f8..3311ce4 100644
--- a/src/tint/lang/wgsl/ast/transform/single_entry_point.cc
+++ b/src/tint/lang/wgsl/ast/transform/single_entry_point.cc
@@ -105,7 +105,7 @@
                         // so that its allocated ID so that it won't be affected by other
                         // stripped away overrides
                         auto* global = sem.Get(override);
-                        const auto* id = b.Id(global->OverrideId());
+                        const auto* id = b.Id(global->Attributes().override_id.value());
                         ctx.InsertFront(override->attributes, id);
                     }
                     b.AST().AddGlobalVariable(ctx.Clone(override));
diff --git a/src/tint/lang/wgsl/ast/transform/substitute_override.cc b/src/tint/lang/wgsl/ast/transform/substitute_override.cc
index 75dd9df..a36bc65 100644
--- a/src/tint/lang/wgsl/ast/transform/substitute_override.cc
+++ b/src/tint/lang/wgsl/ast/transform/substitute_override.cc
@@ -87,7 +87,7 @@
         Type ty = w->type ? ctx.Clone(w->type) : Type{};
 
         // No replacement provided, just clone the override node as a const.
-        auto iter = data->map.find(sem->OverrideId());
+        auto iter = data->map.find(sem->Attributes().override_id.value());
         if (iter == data->map.end()) {
             if (!w->initializer) {
                 b.Diagnostics().add_error(
diff --git a/src/tint/lang/wgsl/ast/transform/vertex_pulling.cc b/src/tint/lang/wgsl/ast/transform/vertex_pulling.cc
index 145d053..0ebf6fa 100644
--- a/src/tint/lang/wgsl/ast/transform/vertex_pulling.cc
+++ b/src/tint/lang/wgsl/ast/transform/vertex_pulling.cc
@@ -796,11 +796,11 @@
             auto* sem = src.Sem().Get<sem::Parameter>(param);
             info.type = sem->Type();
 
-            if (TINT_UNLIKELY(!sem->Location().has_value())) {
+            if (TINT_UNLIKELY(!sem->Attributes().location.has_value())) {
                 TINT_ICE() << "Location missing value";
                 return;
             }
-            location_info[sem->Location().value()] = info;
+            location_info[sem->Attributes().location.value()] = info;
         } else {
             auto* builtin_attr = GetAttribute<BuiltinAttribute>(param->attributes);
             if (TINT_UNLIKELY(!builtin_attr)) {
diff --git a/src/tint/lang/wgsl/extension.cc b/src/tint/lang/wgsl/extension.cc
index a9f363c..e30d7b2 100644
--- a/src/tint/lang/wgsl/extension.cc
+++ b/src/tint/lang/wgsl/extension.cc
@@ -48,6 +48,9 @@
     if (str == "chromium_experimental_dp4a") {
         return Extension::kChromiumExperimentalDp4A;
     }
+    if (str == "chromium_experimental_framebuffer_fetch") {
+        return Extension::kChromiumExperimentalFramebufferFetch;
+    }
     if (str == "chromium_experimental_full_ptr_parameters") {
         return Extension::kChromiumExperimentalFullPtrParameters;
     }
@@ -83,6 +86,8 @@
             return "chromium_disable_uniformity_analysis";
         case Extension::kChromiumExperimentalDp4A:
             return "chromium_experimental_dp4a";
+        case Extension::kChromiumExperimentalFramebufferFetch:
+            return "chromium_experimental_framebuffer_fetch";
         case Extension::kChromiumExperimentalFullPtrParameters:
             return "chromium_experimental_full_ptr_parameters";
         case Extension::kChromiumExperimentalPixelLocal:
diff --git a/src/tint/lang/wgsl/extension.h b/src/tint/lang/wgsl/extension.h
index 931c1ad..0e259f1 100644
--- a/src/tint/lang/wgsl/extension.h
+++ b/src/tint/lang/wgsl/extension.h
@@ -48,6 +48,7 @@
     kUndefined,
     kChromiumDisableUniformityAnalysis,
     kChromiumExperimentalDp4A,
+    kChromiumExperimentalFramebufferFetch,
     kChromiumExperimentalFullPtrParameters,
     kChromiumExperimentalPixelLocal,
     kChromiumExperimentalPushConstant,
@@ -76,17 +77,24 @@
 Extension ParseExtension(std::string_view str);
 
 constexpr std::string_view kExtensionStrings[] = {
-    "chromium_disable_uniformity_analysis",      "chromium_experimental_dp4a",
-    "chromium_experimental_full_ptr_parameters", "chromium_experimental_pixel_local",
-    "chromium_experimental_push_constant",       "chromium_experimental_read_write_storage_texture",
-    "chromium_experimental_subgroups",           "chromium_internal_dual_source_blending",
-    "chromium_internal_relaxed_uniform_layout",  "f16",
+    "chromium_disable_uniformity_analysis",
+    "chromium_experimental_dp4a",
+    "chromium_experimental_framebuffer_fetch",
+    "chromium_experimental_full_ptr_parameters",
+    "chromium_experimental_pixel_local",
+    "chromium_experimental_push_constant",
+    "chromium_experimental_read_write_storage_texture",
+    "chromium_experimental_subgroups",
+    "chromium_internal_dual_source_blending",
+    "chromium_internal_relaxed_uniform_layout",
+    "f16",
 };
 
 /// All extensions
 static constexpr Extension kAllExtensions[] = {
     Extension::kChromiumDisableUniformityAnalysis,
     Extension::kChromiumExperimentalDp4A,
+    Extension::kChromiumExperimentalFramebufferFetch,
     Extension::kChromiumExperimentalFullPtrParameters,
     Extension::kChromiumExperimentalPixelLocal,
     Extension::kChromiumExperimentalPushConstant,
diff --git a/src/tint/lang/wgsl/extension_bench.cc b/src/tint/lang/wgsl/extension_bench.cc
index 5d7a6d9..06b7f67 100644
--- a/src/tint/lang/wgsl/extension_bench.cc
+++ b/src/tint/lang/wgsl/extension_bench.cc
@@ -59,62 +59,69 @@
         "chromium_exverimentiil_dp4a",
         "chro8ium_experimenWWal_dp4a",
         "chromiMm_eperimxxntal_dp4a",
-        "chromium_expeggimeXtal_full_ptr_paraeters",
-        "chromium_expVrimental_full_ptr_puraXeer",
-        "chromium_experimental_full_ptr3parameters",
+        "cXromium_expermental_framggbuffer_fetch",
+        "chVomiu_experimntal_fXauebuffer_fetch",
+        "chromium_experimental_fr3mebuffer_fetch",
+        "chromium_experimental_framebuffer_fetch",
+        "chromium_experimental_fraEebuffer_fetch",
+        "chromTTum_experimental_fraebuffePP_fetch",
+        "chromum_experiddental_framebxxffer_fetch",
+        "chromium_experimental_full_ptr_p44rameters",
+        "SShromium_experimental_full_ptr_parameVVers",
+        "chroRium_expRrimental_fu22l_ptr_prameters",
         "chromium_experimental_full_ptr_parameters",
-        "chromium_experimentalEfull_ptr_parameters",
-        "chromium_experimentalfull_ptr_PPaTTameters",
-        "chromium_ddxperimental_fullptrxxparameters",
-        "chromium_experi44ental_pixel_local",
-        "chromium_experimental_VVSixel_local",
-        "chroRium_experimental_pix22Rlocal",
+        "chromium_experimFntal_full_ptr_paramet9r",
+        "chromium_experimentl_full_ptr_parameters",
+        "chromiuHexperimental_fulOO_pVr_paramRRters",
+        "chromium_experimenya_pixel_local",
+        "chrromium_exp77rimentnnl_pixellllocGl",
+        "chromium_exper4mental_pixel_lo00al",
         "chromium_experimental_pixel_local",
-        "chromiuF_experiment9lpixel_local",
-        "chromium_experimental_pixel_loca",
-        "Vhromium_expeOOimentalHpixRRl_lcal",
-        "chromiym_experimental_push_contant",
-        "nnhro77ium_experimenGal_push_conrrllant",
-        "chromium_experimental_push_c4nstan00",
+        "chooomium_eperimenal_pxel_local",
+        "chromium_experzzental_pixel_ocal",
+        "chiiomiu_epperimental_pix11l_local",
+        "chromXXum_experimental_push_constant",
+        "chromII9um_experinnental_push_c55nstant",
+        "chSSomium_experiHHental_pusrraaconstaYt",
         "chromium_experimental_push_constant",
-        "chooomum_experimental_ush_constat",
-        "chromium_xperimntal_zzush_constant",
-        "chromi11m_experimepptal_psh_ciistant",
-        "chromium_experimental_read_writeXXstorage_texture",
-        "chromium_exII55rimental_read_write99storagnn_texture",
-        "chromiumSSexperaamental_read_wYitrr_storage_HHexture",
+        "chromium_epHrimentkk_psh_constant",
+        "chromium_expegimenja_puRRh_costant",
+        "chromium_bxperimental_push_contan",
+        "chromium_experimental_read_wjite_storage_texture",
+        "chromium_expeimental_read_write_storage_texture",
+        "chqomium_experimentl_rad_write_storage_texture",
         "chromium_experimental_read_write_storage_texture",
-        "Hhromium_experimeta_rkkad_write_strage_texture",
-        "chromium_experijental_red_wrRgte_storag_texture",
-        "chromium_exbeimentalread_write_storage_texture",
-        "chromium_experimental_sjbgroups",
-        "chromium_experimental_sbgroups",
-        "cromum_experimentalqsubgroups",
+        "chromium_experimental_read_writNN_storage_texure",
+        "chromiumexperimental_read_write_stovvage_texure",
+        "chromium_xperimental_read_write_storage_textQQre",
+        "cffromium_experimenralsubgrous",
+        "chromium_experimentjl_subgroups",
+        "chromiu2ww8NNperimental_subgoups",
         "chromium_experimental_subgroups",
-        "chromium_expNNrimental_subgoups",
-        "chromium_experimetal_svvbgrous",
-        "chromium_experiQental_subgroups",
-        "chrorum_internal_dal_source_bleffding",
-        "chromium_internal_dual_source_jlending",
-        "chromiNNm_internal_dua8_sourwwe_blening",
+        "chromium_experimental_subgroup",
+        "chromiurr_experimental_subgroups",
+        "cGromium_experimental_subgroups",
+        "chromium_internFFl_dual_source_blending",
+        "cEromum_internl_dual_source_bleding",
+        "chromium_internal_dual_source_brrendin",
         "chromium_internal_dual_source_blending",
-        "chromium_internal_dual_soure_blending",
-        "chromium_irrternal_dual_source_blending",
-        "chromium_internal_duaG_source_blending",
-        "chromium_internalFFrelaxed_uniform_layout",
-        "chromEum_internal_relaxed_unifrmlyout",
-        "chromium_internalrrrelaxd_uniform_layout",
+        "chromium_internal_dual_surce_blendin",
+        "chromium_iXterJJa_dual_souDce_blending",
+        "chromi8m_ineral_dual_source_blendin",
+        "chromu11_internal_relaed_uniform_kayou",
+        "chromium_internal_relaxd_uniform_layout",
+        "chromium_internal_elaxed_uJiform_layout",
         "chromium_internal_relaxed_uniform_layout",
-        "chromiuminternal_relaxed_uniform_layut",
-        "cXroDium_internal_rJJlaed_uniform_layout",
-        "chromium_int8nal_relaed_uniform_layut",
-        "k",
-        "16",
-        "J1",
+        "chromium_intcrnal_relaxed_uniform_layout",
+        "chromiOm_internal_relaxed_uniform_layout",
+        "chromiutt_intervva_KK_relaxed_uniform_layout",
+        "xx8",
+        "__F",
+        "f1q",
         "f16",
-        "c16",
-        "fO6",
-        "_KKttvv",
+        "331O",
+        "ftt6QQ",
+        "666",
     };
     for (auto _ : state) {
         for (auto* str : kStrings) {
diff --git a/src/tint/lang/wgsl/extension_test.cc b/src/tint/lang/wgsl/extension_test.cc
index 648baea..029deb2 100644
--- a/src/tint/lang/wgsl/extension_test.cc
+++ b/src/tint/lang/wgsl/extension_test.cc
@@ -59,6 +59,7 @@
 static constexpr Case kValidCases[] = {
     {"chromium_disable_uniformity_analysis", Extension::kChromiumDisableUniformityAnalysis},
     {"chromium_experimental_dp4a", Extension::kChromiumExperimentalDp4A},
+    {"chromium_experimental_framebuffer_fetch", Extension::kChromiumExperimentalFramebufferFetch},
     {"chromium_experimental_full_ptr_parameters",
      Extension::kChromiumExperimentalFullPtrParameters},
     {"chromium_experimental_pixel_local", Extension::kChromiumExperimentalPixelLocal},
@@ -78,30 +79,33 @@
     {"chro1ium_experimental_dp4a", Extension::kUndefined},
     {"chrJmium_experiqqetal_dp4a", Extension::kUndefined},
     {"chromium_experimenll77l_dp4a", Extension::kUndefined},
-    {"chroium_experimental_full_ptr_paqqppmetHHrs", Extension::kUndefined},
-    {"chrium_evperiental_full_ptr_paraceters", Extension::kUndefined},
-    {"chromium_expGimental_fullbptr_parameters", Extension::kUndefined},
-    {"vhromium_experimental_pixel_liical", Extension::kUndefined},
-    {"chromium_experiment8l_pixel_lWWcal", Extension::kUndefined},
-    {"chromium_expeimentMl_xxixel_local", Extension::kUndefined},
-    {"chrXmium_experimeggtal_ush_constant", Extension::kUndefined},
-    {"chromiu_experVmentalpusX_constant", Extension::kUndefined},
-    {"chro3ium_experimental_push_constant", Extension::kUndefined},
-    {"chromium_experimentEl_read_write_storage_texture", Extension::kUndefined},
-    {"chromium_experimePPtTT_read_write_storage_texture", Extension::kUndefined},
-    {"chromium_expeimental_read_write_stoddagexxtexture", Extension::kUndefined},
-    {"chromium_experimental_44ubgroups", Extension::kUndefined},
-    {"cSSromVVum_experimental_subgroups", Extension::kUndefined},
-    {"chrmium_e22perimental_suRgrRups", Extension::kUndefined},
-    {"chroFium_internal_dual_source_bl9ndig", Extension::kUndefined},
-    {"chrmium_internal_dual_source_blending", Extension::kUndefined},
-    {"cVromium_interHal_dualOOsouRRce_blening", Extension::kUndefined},
-    {"chromium_internl_relaxyd_uniform_layout", Extension::kUndefined},
-    {"chromnnum_internrr77_Gelaxell_uniform_layout", Extension::kUndefined},
-    {"chromium_intern4l_relaxe00_uniform_layout", Extension::kUndefined},
-    {"5", Extension::kUndefined},
-    {"u16", Extension::kUndefined},
-    {"f", Extension::kUndefined},
+    {"cqqromium_eppperimental_framebuffe_fetcHH", Extension::kUndefined},
+    {"chrmium_experimvntal_frcmebufer_ftch", Extension::kUndefined},
+    {"chromium_expebimental_framGbufer_fetch", Extension::kUndefined},
+    {"chvomium_exiierimental_full_ptr_parameters", Extension::kUndefined},
+    {"chromium_WWxperimental_full8ptr_parameters", Extension::kUndefined},
+    {"chromxxum_Mperimental_full_ptr_parameters", Extension::kUndefined},
+    {"chromum_experimental_pixeX_loggal", Extension::kUndefined},
+    {"chromium_expVrXmntal_ixel_local", Extension::kUndefined},
+    {"3hromium_experimental_pixel_local", Extension::kUndefined},
+    {"chromium_experEmental_push_constant", Extension::kUndefined},
+    {"chPPomiumexperimental_push_conTTtant", Extension::kUndefined},
+    {"chromixxm_experimentddl_push_constnt", Extension::kUndefined},
+    {"44hromium_experimental_read_write_storage_texture", Extension::kUndefined},
+    {"chromium_experimental_reaSS_wriVVe_storage_texture", Extension::kUndefined},
+    {"chro22ium_eRperimental_read_Rrite_storag_texture", Extension::kUndefined},
+    {"chromium_experimental_sbgroup9", Extension::kUndefined},
+    {"cromium_experimental_subgroups", Extension::kUndefined},
+    {"VhrHium_experimental_subOOrouRRs", Extension::kUndefined},
+    {"chromium_internay_dual_sorce_blending", Extension::kUndefined},
+    {"chrnnmium_internal_duGrr_source_bllend77ng", Extension::kUndefined},
+    {"chromiu4_inter00al_dual_source_blending", Extension::kUndefined},
+    {"chrmoom_internal_relaxed_uniform_lyout", Extension::kUndefined},
+    {"chroium_internal_rlaxed_uniform_layzzut", Extension::kUndefined},
+    {"chromium_internaii_r11axed_uppifor_layout", Extension::kUndefined},
+    {"f1XX", Extension::kUndefined},
+    {"55199II", Extension::kUndefined},
+    {"frSSHHa", Extension::kUndefined},
 };
 
 using ExtensionParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/wgsl/helpers/flatten_bindings_test.cc b/src/tint/lang/wgsl/helpers/flatten_bindings_test.cc
index b395664..003ef72 100644
--- a/src/tint/lang/wgsl/helpers/flatten_bindings_test.cc
+++ b/src/tint/lang/wgsl/helpers/flatten_bindings_test.cc
@@ -81,18 +81,18 @@
 
     auto* sem = flattened->Sem().Get<sem::GlobalVariable>(vars[0]);
     ASSERT_NE(sem, nullptr);
-    EXPECT_EQ(sem->BindingPoint()->group, 0u);
-    EXPECT_EQ(sem->BindingPoint()->binding, 0u);
+    EXPECT_EQ(sem->Attributes().binding_point->group, 0u);
+    EXPECT_EQ(sem->Attributes().binding_point->binding, 0u);
 
     sem = flattened->Sem().Get<sem::GlobalVariable>(vars[1]);
     ASSERT_NE(sem, nullptr);
-    EXPECT_EQ(sem->BindingPoint()->group, 0u);
-    EXPECT_EQ(sem->BindingPoint()->binding, 1u);
+    EXPECT_EQ(sem->Attributes().binding_point->group, 0u);
+    EXPECT_EQ(sem->Attributes().binding_point->binding, 1u);
 
     sem = flattened->Sem().Get<sem::GlobalVariable>(vars[2]);
     ASSERT_NE(sem, nullptr);
-    EXPECT_EQ(sem->BindingPoint()->group, 0u);
-    EXPECT_EQ(sem->BindingPoint()->binding, 2u);
+    EXPECT_EQ(sem->Attributes().binding_point->group, 0u);
+    EXPECT_EQ(sem->Attributes().binding_point->binding, 2u);
 }
 
 TEST_F(FlattenBindingsTest, NotFlat_MultipleNamespaces) {
@@ -144,20 +144,20 @@
     for (size_t i = 0; i < num_buffers; ++i) {
         auto* sem = flattened->Sem().Get<sem::GlobalVariable>(vars[i]);
         ASSERT_NE(sem, nullptr);
-        EXPECT_EQ(sem->BindingPoint()->group, 0u);
-        EXPECT_EQ(sem->BindingPoint()->binding, i);
+        EXPECT_EQ(sem->Attributes().binding_point->group, 0u);
+        EXPECT_EQ(sem->Attributes().binding_point->binding, i);
     }
     for (size_t i = 0; i < num_samplers; ++i) {
         auto* sem = flattened->Sem().Get<sem::GlobalVariable>(vars[i + num_buffers]);
         ASSERT_NE(sem, nullptr);
-        EXPECT_EQ(sem->BindingPoint()->group, 0u);
-        EXPECT_EQ(sem->BindingPoint()->binding, i);
+        EXPECT_EQ(sem->Attributes().binding_point->group, 0u);
+        EXPECT_EQ(sem->Attributes().binding_point->binding, i);
     }
     for (size_t i = 0; i < num_textures; ++i) {
         auto* sem = flattened->Sem().Get<sem::GlobalVariable>(vars[i + num_buffers + num_samplers]);
         ASSERT_NE(sem, nullptr);
-        EXPECT_EQ(sem->BindingPoint()->group, 0u);
-        EXPECT_EQ(sem->BindingPoint()->binding, i);
+        EXPECT_EQ(sem->Attributes().binding_point->group, 0u);
+        EXPECT_EQ(sem->Attributes().binding_point->binding, i);
     }
 }
 
diff --git a/src/tint/lang/wgsl/inspector/inspector.cc b/src/tint/lang/wgsl/inspector/inspector.cc
index 278773a..7f49f66 100644
--- a/src/tint/lang/wgsl/inspector/inspector.cc
+++ b/src/tint/lang/wgsl/inspector/inspector.cc
@@ -177,7 +177,7 @@
     for (auto* param : sem->Parameters()) {
         AddEntryPointInOutVariables(param->Declaration()->name->symbol.Name(),
                                     param->Declaration()->name->symbol.Name(), param->Type(),
-                                    param->Declaration()->attributes, param->Location(),
+                                    param->Declaration()->attributes, param->Attributes().location,
                                     entry_point.input_variables);
 
         entry_point.input_position_used |= ContainsBuiltin(
@@ -212,10 +212,10 @@
         auto name = decl->name->symbol.Name();
 
         auto* global = var->As<sem::GlobalVariable>();
-        if (global && global->Declaration()->Is<ast::Override>()) {
+        if (auto override_id = global->Attributes().override_id) {
             Override override;
             override.name = name;
-            override.id = global->OverrideId();
+            override.id = override_id.value();
             auto* type = var->Type();
             TINT_ASSERT(type->Is<core::type::Scalar>());
             if (type->is_bool_scalar_or_vector()) {
@@ -279,7 +279,7 @@
         // WGSL, so the resolver should catch it. Thus here the inspector just
         // assumes all definitions of the override id are the same, so only needs
         // to find the first reference to override id.
-        OverrideId override_id = global->OverrideId();
+        auto override_id = global->Attributes().override_id.value();
         if (result.find(override_id) != result.end()) {
             continue;
         }
@@ -311,9 +311,9 @@
     std::map<std::string, OverrideId> result;
     for (auto* var : program_.AST().GlobalVariables()) {
         auto* global = program_.Sem().Get<sem::GlobalVariable>(var);
-        if (global && global->Declaration()->Is<ast::Override>()) {
+        if (auto override_id = global->Attributes().override_id) {
             auto name = var->name->symbol.Name();
-            result[name] = global->OverrideId();
+            result[name] = override_id.value();
         }
     }
     return result;
@@ -527,8 +527,9 @@
         auto* texture = pair.first->As<sem::GlobalVariable>();
         auto* sampler = pair.second ? pair.second->As<sem::GlobalVariable>() : nullptr;
         SamplerTexturePair new_pair;
-        new_pair.sampler_binding_point = sampler ? *sampler->BindingPoint() : placeholder;
-        new_pair.texture_binding_point = *texture->BindingPoint();
+        new_pair.sampler_binding_point =
+            sampler ? *sampler->Attributes().binding_point : placeholder;
+        new_pair.texture_binding_point = *texture->Attributes().binding_point;
         new_pairs.push_back(new_pair);
     }
     return new_pairs;
@@ -818,18 +819,18 @@
         auto* t = c->args[static_cast<size_t>(texture_index)];
         auto* s = c->args[static_cast<size_t>(sampler_index)];
 
-        GetOriginatingResources(std::array<const ast::Expression*, 2>{t, s},
-                                [&](std::array<const sem::GlobalVariable*, 2> globals) {
-                                    auto texture_binding_point = *globals[0]->BindingPoint();
-                                    auto sampler_binding_point = *globals[1]->BindingPoint();
+        GetOriginatingResources(
+            std::array<const ast::Expression*, 2>{t, s},
+            [&](std::array<const sem::GlobalVariable*, 2> globals) {
+                auto texture_binding_point = *globals[0]->Attributes().binding_point;
+                auto sampler_binding_point = *globals[1]->Attributes().binding_point;
 
-                                    for (auto* entry_point : entry_points) {
-                                        const auto& ep_name =
-                                            entry_point->Declaration()->name->symbol.Name();
-                                        (*sampler_targets_)[ep_name].Add(
-                                            {sampler_binding_point, texture_binding_point});
-                                    }
-                                });
+                for (auto* entry_point : entry_points) {
+                    const auto& ep_name = entry_point->Declaration()->name->symbol.Name();
+                    (*sampler_targets_)[ep_name].Add(
+                        {sampler_binding_point, texture_binding_point});
+                }
+            });
     }
 }
 
diff --git a/src/tint/lang/wgsl/inspector/inspector_test.cc b/src/tint/lang/wgsl/inspector/inspector_test.cc
index 052c464..0eb5c8e 100644
--- a/src/tint/lang/wgsl/inspector/inspector_test.cc
+++ b/src/tint/lang/wgsl/inspector/inspector_test.cc
@@ -1879,15 +1879,15 @@
 
     ASSERT_TRUE(result.count("a"));
     ASSERT_TRUE(program_->Sem().Get(a));
-    EXPECT_EQ(result["a"], program_->Sem().Get(a)->OverrideId());
+    EXPECT_EQ(result["a"], program_->Sem().Get(a)->Attributes().override_id);
 
     ASSERT_TRUE(result.count("b"));
     ASSERT_TRUE(program_->Sem().Get(b));
-    EXPECT_EQ(result["b"], program_->Sem().Get(b)->OverrideId());
+    EXPECT_EQ(result["b"], program_->Sem().Get(b)->Attributes().override_id);
 
     ASSERT_TRUE(result.count("c"));
     ASSERT_TRUE(program_->Sem().Get(c));
-    EXPECT_EQ(result["c"], program_->Sem().Get(c)->OverrideId());
+    EXPECT_EQ(result["c"], program_->Sem().Get(c)->Attributes().override_id);
 }
 
 TEST_F(InspectorGetResourceBindingsTest, Empty) {
diff --git a/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc b/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc
index 26bd8dc..893573c 100644
--- a/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc
@@ -205,7 +205,7 @@
     // Error when unknown extension found
     EXPECT_TRUE(p->has_error());
     EXPECT_EQ(p->error(), R"(1:8: expected extension
-Possible values: 'chromium_disable_uniformity_analysis', 'chromium_experimental_dp4a', 'chromium_experimental_full_ptr_parameters', 'chromium_experimental_pixel_local', 'chromium_experimental_push_constant', 'chromium_experimental_read_write_storage_texture', 'chromium_experimental_subgroups', 'chromium_internal_dual_source_blending', 'chromium_internal_relaxed_uniform_layout', 'f16')");
+Possible values: 'chromium_disable_uniformity_analysis', 'chromium_experimental_dp4a', 'chromium_experimental_framebuffer_fetch', 'chromium_experimental_full_ptr_parameters', 'chromium_experimental_pixel_local', 'chromium_experimental_push_constant', 'chromium_experimental_read_write_storage_texture', 'chromium_experimental_subgroups', 'chromium_internal_dual_source_blending', 'chromium_internal_relaxed_uniform_layout', 'f16')");
     auto program = p->program();
     auto& ast = program.AST();
     EXPECT_EQ(ast.Enables().Length(), 0u);
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
index a0c10a8..5e99225 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
@@ -456,12 +456,8 @@
                         }
                     });
 
-                if (param_sem->Location().has_value()) {
-                    param->SetLocation(param_sem->Location().value(), interpolation);
-                }
-                if (param_sem->BindingPoint().has_value()) {
-                    param->SetBindingPoint(param_sem->BindingPoint()->group,
-                                           param_sem->BindingPoint()->binding);
+                if (param_sem->Attributes().location.has_value()) {
+                    param->SetLocation(param_sem->Attributes().location.value(), interpolation);
                 }
             }
 
@@ -1324,8 +1320,8 @@
                 current_block_->Append(val);
 
                 if (auto* gv = sem->As<sem::GlobalVariable>(); gv && var->HasBindingPoint()) {
-                    val->SetBindingPoint(gv->BindingPoint().value().group,
-                                         gv->BindingPoint().value().binding);
+                    val->SetBindingPoint(gv->Attributes().binding_point->group,
+                                         gv->Attributes().binding_point->binding);
                 }
 
                 // Store the declaration so we can get the instruction to store too
diff --git a/src/tint/lang/wgsl/resolver/attribute_validation_test.cc b/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
index 64e2bc7..d88daa5 100644
--- a/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
@@ -31,6 +31,8 @@
 #include "src/tint/lang/wgsl/ast/transform/add_block_attribute.h"
 #include "src/tint/lang/wgsl/resolver/resolver.h"
 #include "src/tint/lang/wgsl/resolver/resolver_helper_test.h"
+#include "src/tint/utils/containers/transform.h"
+#include "src/tint/utils/macros/compiler.h"
 #include "src/tint/utils/text/string_stream.h"
 
 #include "gmock/gmock.h"
@@ -57,7 +59,7 @@
 enum class AttributeKind {
     kAlign,
     kBinding,
-    kBuiltin,
+    kBuiltinPosition,
     kDiagnostic,
     kGroup,
     kId,
@@ -68,18 +70,53 @@
     kMustUse,
     kOffset,
     kSize,
-    kStage,
+    kStageCompute,
     kStride,
-    kWorkgroup,
-
-    kBindingAndGroup,
+    kWorkgroupSize,
 };
+static std::ostream& operator<<(std::ostream& o, AttributeKind k) {
+    switch (k) {
+        case AttributeKind::kAlign:
+            return o << "@align";
+        case AttributeKind::kBinding:
+            return o << "@binding";
+        case AttributeKind::kBuiltinPosition:
+            return o << "@builtin(position)";
+        case AttributeKind::kDiagnostic:
+            return o << "@diagnostic";
+        case AttributeKind::kGroup:
+            return o << "@group";
+        case AttributeKind::kId:
+            return o << "@id";
+        case AttributeKind::kIndex:
+            return o << "@index";
+        case AttributeKind::kInterpolate:
+            return o << "@interpolate";
+        case AttributeKind::kInvariant:
+            return o << "@invariant";
+        case AttributeKind::kLocation:
+            return o << "@location";
+        case AttributeKind::kOffset:
+            return o << "@offset";
+        case AttributeKind::kMustUse:
+            return o << "@must_use";
+        case AttributeKind::kSize:
+            return o << "@size";
+        case AttributeKind::kStageCompute:
+            return o << "@stage(compute)";
+        case AttributeKind::kStride:
+            return o << "@stride";
+        case AttributeKind::kWorkgroupSize:
+            return o << "@workgroup_size";
+    }
+    TINT_UNREACHABLE();
+    return o << "<unknown>";
+}
 
 static bool IsBindingAttribute(AttributeKind kind) {
     switch (kind) {
         case AttributeKind::kBinding:
         case AttributeKind::kGroup:
-        case AttributeKind::kBindingAndGroup:
             return true;
         default:
             return false;
@@ -87,195 +124,500 @@
 }
 
 struct TestParams {
-    AttributeKind kind;
-    bool should_pass;
+    Vector<AttributeKind, 2> attributes;
+    std::string error;  // empty string (Pass) is an expected pass
 };
+
+static constexpr const char* Pass = "";
+
+static std::vector<TestParams> OnlyDiagnosticValidFor(std::string thing) {
+    return {TestParams{
+                {AttributeKind::kAlign},
+                "1:2 error: @align is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kBinding},
+                "1:2 error: @binding is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kBuiltinPosition},
+                "1:2 error: @builtin is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kDiagnostic},
+                Pass,
+            },
+            TestParams{
+                {AttributeKind::kGroup},
+                "1:2 error: @group is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kId},
+                "1:2 error: @id is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kIndex},
+                "1:2 error: @index is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kInterpolate},
+                "1:2 error: @interpolate is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kInvariant},
+                "1:2 error: @invariant is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kLocation},
+                "1:2 error: @location is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kMustUse},
+                "1:2 error: @must_use is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kOffset},
+                "1:2 error: @offset is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kSize},
+                "1:2 error: @size is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kStageCompute},
+                "1:2 error: @stage is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kStride},
+                "1:2 error: @stride is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kWorkgroupSize},
+                "1:2 error: @workgroup_size is not valid for " + thing,
+            },
+            TestParams{
+                {AttributeKind::kBinding, AttributeKind::kGroup},
+                "1:2 error: @binding is not valid for " + thing,
+            }};
+}
+
+static std::ostream& operator<<(std::ostream& o, const TestParams& c) {
+    return o << "attributes: " << c.attributes << ", expect pass: " << c.error.empty();
+}
+
+const ast::Attribute* CreateAttribute(const Source& source,
+                                      ProgramBuilder& builder,
+                                      AttributeKind kind) {
+    switch (kind) {
+        case AttributeKind::kAlign:
+            return builder.MemberAlign(source, 4_i);
+        case AttributeKind::kBinding:
+            return builder.Binding(source, 1_a);
+        case AttributeKind::kBuiltinPosition:
+            return builder.Builtin(source, core::BuiltinValue::kPosition);
+        case AttributeKind::kDiagnostic:
+            return builder.DiagnosticAttribute(source, wgsl::DiagnosticSeverity::kInfo, "chromium",
+                                               "unreachable_code");
+        case AttributeKind::kGroup:
+            return builder.Group(source, 1_a);
+        case AttributeKind::kId:
+            return builder.Id(source, 0_a);
+        case AttributeKind::kIndex:
+            return builder.Index(source, 0_a);
+        case AttributeKind::kInterpolate:
+            return builder.Interpolate(source, core::InterpolationType::kLinear,
+                                       core::InterpolationSampling::kCenter);
+        case AttributeKind::kInvariant:
+            return builder.Invariant(source);
+        case AttributeKind::kLocation:
+            return builder.Location(source, 0_a);
+        case AttributeKind::kOffset:
+            return builder.MemberOffset(source, 4_a);
+        case AttributeKind::kMustUse:
+            return builder.MustUse(source);
+        case AttributeKind::kSize:
+            return builder.MemberSize(source, 16_a);
+        case AttributeKind::kStageCompute:
+            return builder.Stage(source, ast::PipelineStage::kCompute);
+        case AttributeKind::kStride:
+            return builder.create<ast::StrideAttribute>(source, 4u);
+        case AttributeKind::kWorkgroupSize:
+            return builder.create<ast::WorkgroupAttribute>(source, builder.Expr(1_i));
+    }
+    TINT_UNREACHABLE() << kind;
+    return nullptr;
+}
+
 struct TestWithParams : ResolverTestWithParam<TestParams> {
-    void EnableExtensionIfNecessary(AttributeKind attributeKind) {
-        if (attributeKind == AttributeKind::kIndex) {
+    void EnableExtensionIfNecessary(AttributeKind attribute) {
+        if (attribute == AttributeKind::kIndex) {
             Enable(wgsl::Extension::kChromiumInternalDualSourceBlending);
         }
     }
+
+    void EnableRequiredExtensions() {
+        for (auto attribute : GetParam().attributes) {
+            EnableExtensionIfNecessary(attribute);
+        }
+    }
+
+    Vector<const ast::Attribute*, 2> CreateAttributes(ProgramBuilder& builder,
+                                                      VectorRef<AttributeKind> kinds) {
+        return Transform<2>(kinds, [&](AttributeKind kind, size_t index) {
+            return CreateAttribute(Source{{static_cast<uint32_t>(index) * 2 + 1,
+                                           static_cast<uint32_t>(index) * 2 + 2}},
+                                   builder, kind);
+        });
+    }
+
+    Vector<const ast::Attribute*, 2> CreateAttributes() {
+        return CreateAttributes(*this, GetParam().attributes);
+    }
 };
 
-static Vector<const ast::Attribute*, 2> createAttributes(const Source& source,
-                                                         ProgramBuilder& builder,
-                                                         AttributeKind kind) {
-    switch (kind) {
-        case AttributeKind::kAlign:
-            return {builder.MemberAlign(source, 4_i)};
-        case AttributeKind::kBinding:
-            return {builder.Binding(source, 1_a)};
-        case AttributeKind::kBuiltin:
-            return {builder.Builtin(source, core::BuiltinValue::kPosition)};
-        case AttributeKind::kDiagnostic:
-            return {builder.DiagnosticAttribute(source, wgsl::DiagnosticSeverity::kInfo, "chromium",
-                                                "unreachable_code")};
-        case AttributeKind::kGroup:
-            return {builder.Group(source, 1_a)};
-        case AttributeKind::kId:
-            return {builder.Id(source, 0_a)};
-        case AttributeKind::kIndex:
-            return {builder.Index(source, 0_a)};
-        case AttributeKind::kInterpolate:
-            return {builder.Interpolate(source, core::InterpolationType::kLinear,
-                                        core::InterpolationSampling::kCenter)};
-        case AttributeKind::kInvariant:
-            return {builder.Invariant(source)};
-        case AttributeKind::kLocation:
-            return {builder.Location(source, 1_a)};
-        case AttributeKind::kOffset:
-            return {builder.MemberOffset(source, 4_a)};
-        case AttributeKind::kMustUse:
-            return {builder.MustUse(source)};
-        case AttributeKind::kSize:
-            return {builder.MemberSize(source, 16_a)};
-        case AttributeKind::kStage:
-            return {builder.Stage(source, ast::PipelineStage::kCompute)};
-        case AttributeKind::kStride:
-            return {builder.create<ast::StrideAttribute>(source, 4u)};
-        case AttributeKind::kWorkgroup:
-            return {builder.create<ast::WorkgroupAttribute>(source, builder.Expr(1_i))};
-        case AttributeKind::kBindingAndGroup:
-            return {builder.Binding(source, 1_a), builder.Group(source, 1_a)};
-    }
-    return {};
-}
+#define CHECK()                                      \
+    if (GetParam().error.empty()) {                  \
+        EXPECT_TRUE(r()->Resolve()) << r()->error(); \
+    } else {                                         \
+        EXPECT_FALSE(r()->Resolve());                \
+        EXPECT_EQ(GetParam().error, r()->error());   \
+    }                                                \
+    TINT_REQUIRE_SEMICOLON
 
-static std::string name(AttributeKind kind) {
-    switch (kind) {
-        case AttributeKind::kAlign:
-            return "@align";
-        case AttributeKind::kBinding:
-            return "@binding";
-        case AttributeKind::kBuiltin:
-            return "@builtin";
-        case AttributeKind::kDiagnostic:
-            return "@diagnostic";
-        case AttributeKind::kGroup:
-            return "@group";
-        case AttributeKind::kId:
-            return "@id";
-        case AttributeKind::kIndex:
-            return "@index";
-        case AttributeKind::kInterpolate:
-            return "@interpolate";
-        case AttributeKind::kInvariant:
-            return "@invariant";
-        case AttributeKind::kLocation:
-            return "@location";
-        case AttributeKind::kOffset:
-            return "@offset";
-        case AttributeKind::kMustUse:
-            return "@must_use";
-        case AttributeKind::kSize:
-            return "@size";
-        case AttributeKind::kStage:
-            return "@stage";
-        case AttributeKind::kStride:
-            return "@stride";
-        case AttributeKind::kWorkgroup:
-            return "@workgroup_size";
-        case AttributeKind::kBindingAndGroup:
-            return "@binding";
-    }
-    return "<unknown>";
+namespace FunctionTests {
+using VoidFunctionAttributeTest = TestWithParams;
+TEST_P(VoidFunctionAttributeTest, IsValid) {
+    EnableRequiredExtensions();
+
+    Func(Source{{9, 9}}, "main", Empty, ty.void_(), Empty, CreateAttributes());
+
+    CHECK();
 }
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    VoidFunctionAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: @binding is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: @builtin is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: @group is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: @id is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(1:2 error: @index is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: @interpolate is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: @invariant is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: @location is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use can only be applied to functions that return a value)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(9:9 error: a compute shader must include 'workgroup_size' in its attributes)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute, AttributeKind::kWorkgroupSize},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for functions)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is only valid for compute stages)",
+        }));
+
+using NonVoidFunctionAttributeTest = TestWithParams;
+TEST_P(NonVoidFunctionAttributeTest, IsValid) {
+    EnableRequiredExtensions();
+
+    Func(Source{{9, 9}}, "main", Empty, ty.i32(), Vector{Return(1_i)}, CreateAttributes());
+
+    CHECK();
+}
+INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
+                         NonVoidFunctionAttributeTest,
+                         testing::Values(
+                             TestParams{
+                                 {AttributeKind::kAlign},
+                                 R"(1:2 error: @align is not valid for functions)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kBinding},
+                                 R"(1:2 error: @binding is not valid for functions)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kBuiltinPosition},
+                                 R"(1:2 error: @builtin is not valid for functions)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kDiagnostic},
+                                 Pass,
+                             },
+                             TestParams{
+                                 {AttributeKind::kGroup},
+                                 R"(1:2 error: @group is not valid for functions)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kId},
+                                 R"(1:2 error: @id is not valid for functions)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kIndex},
+                                 R"(1:2 error: @index is not valid for functions)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kInterpolate},
+                                 R"(1:2 error: @interpolate is not valid for functions)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kInvariant},
+                                 R"(1:2 error: @invariant is not valid for functions)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kLocation},
+                                 R"(1:2 error: @location is not valid for functions)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kMustUse},
+                                 Pass,
+                             },
+                             TestParams{
+                                 {AttributeKind::kOffset},
+                                 R"(1:2 error: @offset is not valid for functions)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kSize},
+                                 R"(1:2 error: @size is not valid for functions)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kStageCompute},
+                                 R"(9:9 error: missing entry point IO attribute on return type)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kStageCompute, AttributeKind::kWorkgroupSize},
+                                 R"(9:9 error: missing entry point IO attribute on return type)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kStride},
+                                 R"(1:2 error: @stride is not valid for functions)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kWorkgroupSize},
+                                 R"(1:2 error: @workgroup_size is only valid for compute stages)",
+                             }));
+}  // namespace FunctionTests
+
 namespace FunctionInputAndOutputTests {
 using FunctionParameterAttributeTest = TestWithParams;
 TEST_P(FunctionParameterAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
     Func("main",
          Vector{
-             Param("a", ty.vec4<f32>(), createAttributes({}, *this, params.kind)),
+             Param("a", ty.vec4<f32>(), CreateAttributes()),
          },
          ty.void_(), tint::Empty);
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else if (params.kind == AttributeKind::kLocation || params.kind == AttributeKind::kBuiltin ||
-               params.kind == AttributeKind::kInvariant ||
-               params.kind == AttributeKind::kInterpolate) {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(), "error: " + name(params.kind) +
-                                    " is not valid for non-entry point function parameters");
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "error: " + name(params.kind) + " is not valid for function parameters");
-    }
+    CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         FunctionParameterAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    FunctionParameterAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: @binding is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: @builtin is not valid for non-entry point function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: @diagnostic is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: @group is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: @id is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(1:2 error: @index is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: @interpolate is not valid for non-entry point function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: @invariant is not valid for non-entry point function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: @location is not valid for non-entry point function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: @stage is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is not valid for function parameters)",
+        }));
 
 using FunctionReturnTypeAttributeTest = TestWithParams;
 TEST_P(FunctionReturnTypeAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
     Func("main", tint::Empty, ty.f32(),
          Vector{
              Return(1_f),
          },
-         tint::Empty, createAttributes({}, *this, params.kind));
+         tint::Empty, CreateAttributes());
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(), "error: " + name(params.kind) +
-                                    " is not valid for non-entry point function "
-                                    "return types");
-    }
+    CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         FunctionReturnTypeAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    FunctionReturnTypeAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: @binding is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: @builtin is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: @diagnostic is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: @group is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: @id is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(1:2 error: @index is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: @interpolate is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: @invariant is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: @location is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: @stage is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for non-entry point function return types)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is not valid for non-entry point function return types)",
+        }));
 }  // namespace FunctionInputAndOutputTests
 
 namespace EntryPointInputAndOutputTests {
 using ComputeShaderParameterAttributeTest = TestWithParams;
 TEST_P(ComputeShaderParameterAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
     Func("main",
          Vector{
-             Param("a", ty.vec4<f32>(), createAttributes(Source{{12, 34}}, *this, params.kind)),
+             Param("a", ty.vec4<f32>(), CreateAttributes()),
          },
          ty.void_(), tint::Empty,
          Vector{
@@ -283,97 +625,174 @@
              WorkgroupSize(1_i),
          });
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        if (params.kind == AttributeKind::kBuiltin) {
-            EXPECT_EQ(
-                r()->error(),
-                R"(12:34 error: @builtin(position) cannot be used in input of compute pipeline stage)");
-        } else if (params.kind == AttributeKind::kInterpolate ||
-                   params.kind == AttributeKind::kLocation ||
-                   params.kind == AttributeKind::kInvariant) {
-            EXPECT_EQ(r()->error(), "12:34 error: " + name(params.kind) +
-                                        " is not valid for compute shader inputs");
-        } else {
-            EXPECT_EQ(r()->error(), "12:34 error: " + name(params.kind) +
-                                        " is not valid for function parameters");
-        }
-    }
+    CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         ComputeShaderParameterAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    ComputeShaderParameterAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: @binding is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: @builtin(position) cannot be used for compute shader input)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: @diagnostic is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: @group is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: @id is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(1:2 error: @index is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: @interpolate cannot be used by compute shaders)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: @invariant cannot be used by compute shaders)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: @location cannot be used by compute shaders)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: @stage is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is not valid for function parameters)",
+        }));
 
 using FragmentShaderParameterAttributeTest = TestWithParams;
 TEST_P(FragmentShaderParameterAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
-    auto attrs = createAttributes(Source{{12, 34}}, *this, params.kind);
-    if (params.kind != AttributeKind::kBuiltin && params.kind != AttributeKind::kLocation) {
-        attrs.Push(Builtin(Source{{34, 56}}, core::BuiltinValue::kPosition));
-    }
-    auto* p = Param("a", ty.vec4<f32>(), attrs);
+    EnableRequiredExtensions();
+    auto* p = Param(Source{{9, 9}}, "a", ty.vec4<f32>(), CreateAttributes());
     Func("frag_main", Vector{p}, ty.void_(), tint::Empty,
          Vector{
              Stage(ast::PipelineStage::kFragment),
          });
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: " + name(params.kind) + " is not valid for function parameters");
-    }
+    CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         FragmentShaderParameterAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, true},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         // kInterpolate tested separately (requires @location)
-                                         TestParams{AttributeKind::kInvariant, true},
-                                         TestParams{AttributeKind::kLocation, true},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    FragmentShaderParameterAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: @binding is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: @diagnostic is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: @group is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: @id is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(1:2 error: @index is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(9:9 error: missing entry point IO attribute on parameter)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate, AttributeKind::kBuiltinPosition},
+            R"(1:2 error: @interpolate can only be used with @location)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate, AttributeKind::kLocation},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(9:9 error: missing entry point IO attribute on parameter)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant, AttributeKind::kBuiltinPosition},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: @stage is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is not valid for function parameters)",
+        }));
 
 using VertexShaderParameterAttributeTest = TestWithParams;
 TEST_P(VertexShaderParameterAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    auto attrs = createAttributes(Source{{12, 34}}, *this, params.kind);
-    if (params.kind != AttributeKind::kLocation) {
-        attrs.Push(Location(Source{{34, 56}}, 2_a));
-    }
-    auto* p = Param("a", ty.vec4<f32>(), attrs);
+    auto* p = Param(Source{{9, 9}}, "a", ty.vec4<f32>(), CreateAttributes());
     Func("vertex_main", Vector{p}, ty.vec4<f32>(),
          Vector{
              Return(Call<vec4<f32>>()),
@@ -385,48 +804,96 @@
              Builtin(core::BuiltinValue::kPosition),
          });
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        if (params.kind == AttributeKind::kBuiltin) {
-            EXPECT_EQ(
-                r()->error(),
-                R"(12:34 error: @builtin(position) cannot be used in input of vertex pipeline stage)");
-        } else if (params.kind == AttributeKind::kInvariant) {
-            EXPECT_EQ(r()->error(),
-                      "12:34 error: invariant attribute must only be applied to a "
-                      "position builtin");
-        } else {
-            EXPECT_EQ(r()->error(), "12:34 error: " + name(params.kind) +
-                                        " is not valid for function parameters");
-        }
-    }
+    CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         VertexShaderParameterAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, true},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, true},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    VertexShaderParameterAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: @binding is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: @builtin(position) cannot be used for vertex shader input)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: @diagnostic is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: @group is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: @id is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(1:2 error: @index is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(9:9 error: missing entry point IO attribute on parameter)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate, AttributeKind::kLocation},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kInterpolate, AttributeKind::kBuiltinPosition},
+            R"(3:4 error: @builtin(position) cannot be used for vertex shader input)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(9:9 error: missing entry point IO attribute on parameter)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant, AttributeKind::kLocation},
+            R"(1:2 error: @invariant must be applied to a position builtin)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant, AttributeKind::kBuiltinPosition},
+            R"(3:4 error: @builtin(position) cannot be used for vertex shader input)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: @stage is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for function parameters)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is not valid for function parameters)",
+        }));
 
 using ComputeShaderReturnTypeAttributeTest = TestWithParams;
 TEST_P(ComputeShaderReturnTypeAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
     Func("main", tint::Empty, ty.vec4<f32>(),
          Vector{
@@ -436,112 +903,188 @@
              Stage(ast::PipelineStage::kCompute),
              WorkgroupSize(1_i),
          },
-         createAttributes(Source{{12, 34}}, *this, params.kind));
+         CreateAttributes());
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        if (params.kind == AttributeKind::kBuiltin) {
-            EXPECT_EQ(
-                r()->error(),
-                R"(12:34 error: @builtin(position) cannot be used in output of compute pipeline stage)");
-        } else if (params.kind == AttributeKind::kInterpolate ||
-                   params.kind == AttributeKind::kLocation ||
-                   params.kind == AttributeKind::kInvariant ||
-                   params.kind == AttributeKind::kIndex) {
-            EXPECT_EQ(r()->error(), "12:34 error: " + name(params.kind) +
-                                        " is not valid for compute shader output");
-        } else {
-            EXPECT_EQ(r()->error(), "12:34 error: " + name(params.kind) +
-                                        " is not valid for entry point return "
-                                        "types");
-        }
-    }
+    CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         ComputeShaderReturnTypeAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    ComputeShaderReturnTypeAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: @binding is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: @builtin(position) cannot be used for compute shader output)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: @diagnostic is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: @group is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: @id is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(1:2 error: @index can only be used for fragment shader output)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: @interpolate cannot be used by compute shaders)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: @invariant cannot be used by compute shaders)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: @location cannot be used by compute shaders)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: @stage is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is not valid for entry point return types)",
+        }));
 
 using FragmentShaderReturnTypeAttributeTest = TestWithParams;
 TEST_P(FragmentShaderReturnTypeAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    auto attrs = createAttributes(Source{{12, 34}}, *this, params.kind);
-    attrs.Push(Location(Source{{34, 56}}, 2_a));
-    Func("frag_main", tint::Empty, ty.vec4<f32>(), Vector{Return(Call<vec4<f32>>())},
+    Func(Source{{9, 9}}, "frag_main", tint::Empty, ty.vec4<f32>(),
+         Vector{Return(Call<vec4<f32>>())},
          Vector{
              Stage(ast::PipelineStage::kFragment),
          },
-         attrs);
+         CreateAttributes());
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        if (params.kind == AttributeKind::kBuiltin) {
-            EXPECT_EQ(
-                r()->error(),
-                R"(12:34 error: @builtin(position) cannot be used in output of fragment pipeline stage)");
-        } else if (params.kind == AttributeKind::kInvariant) {
-            EXPECT_EQ(
-                r()->error(),
-                R"(12:34 error: invariant attribute must only be applied to a position builtin)");
-        } else if (params.kind == AttributeKind::kLocation) {
-            EXPECT_EQ(r()->error(),
-                      R"(34:56 error: duplicate location attribute
-12:34 note: first attribute declared here)");
-        } else {
-            EXPECT_EQ(r()->error(), "12:34 error: " + name(params.kind) +
-                                        " is not valid for entry point return types");
-        }
-    }
+    CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         FragmentShaderReturnTypeAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, true},
-                                         TestParams{AttributeKind::kInterpolate, true},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    FragmentShaderReturnTypeAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: @binding is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: @builtin(position) cannot be used for fragment shader output)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: @diagnostic is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: @group is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: @id is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(9:9 error: missing entry point IO attribute on return type)",
+        },
+        TestParams{
+            {AttributeKind::kIndex, AttributeKind::kLocation},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(9:9 error: missing entry point IO attribute on return type)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate, AttributeKind::kLocation},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(9:9 error: missing entry point IO attribute on return type)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant, AttributeKind::kLocation},
+            R"(1:2 error: @invariant must be applied to a position builtin)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: @stage is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kBinding, AttributeKind::kGroup},
+            R"(1:2 error: @binding is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kIndex, AttributeKind::kLocation},
+            Pass,
+        }));
 
 using VertexShaderReturnTypeAttributeTest = TestWithParams;
 TEST_P(VertexShaderReturnTypeAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
-    auto attrs = createAttributes(Source{{12, 34}}, *this, params.kind);
+    EnableRequiredExtensions();
+    auto attrs = CreateAttributes();
     // a vertex shader must include the 'position' builtin in its return type
-    if (params.kind != AttributeKind::kBuiltin) {
-        attrs.Push(Builtin(Source{{34, 56}}, core::BuiltinValue::kPosition));
+    if (!GetParam().attributes.Any([](auto b) { return b == AttributeKind::kBuiltinPosition; })) {
+        attrs.Push(Builtin(Source{{9, 9}}, core::BuiltinValue::kPosition));
     }
     Func("vertex_main", tint::Empty, ty.vec4<f32>(),
          Vector{
@@ -550,64 +1093,90 @@
          Vector{
              Stage(ast::PipelineStage::kVertex),
          },
-         attrs);
+         std::move(attrs));
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        if (params.kind == AttributeKind::kLocation) {
-            EXPECT_EQ(r()->error(),
-                      R"(34:56 error: multiple entry point IO attributes
-12:34 note: previously consumed @location)");
-        } else if (params.kind == AttributeKind::kIndex) {
-            EXPECT_EQ(r()->error(), R"(12:34 error: @index is not valid for vertex shader output)");
-        } else {
-            EXPECT_EQ(r()->error(), "12:34 error: " + name(params.kind) +
-                                        " is not valid for entry point return types");
-        }
-    }
+    CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         VertexShaderReturnTypeAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, true},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         // kInterpolate tested separately (requires @location)
-                                         TestParams{AttributeKind::kInvariant, true},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    VertexShaderReturnTypeAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: @binding is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: @diagnostic is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: @group is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: @id is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(1:2 error: @index can only be used for fragment shader output)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: @interpolate can only be used with @location)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(9:9 error: multiple entry point IO attributes
+1:2 note: previously consumed @location)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: @stage is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kBinding, AttributeKind::kGroup},
+            R"(1:2 error: @binding is not valid for entry point return types)",
+        },
+        TestParams{
+            {AttributeKind::kLocation, AttributeKind::kLocation},
+            R"(3:4 error: duplicate location attribute
+1:2 note: first attribute declared here)",
+        }));
 
 using EntryPointParameterAttributeTest = TestWithParams;
-TEST_F(EntryPointParameterAttributeTest, DuplicateAttribute) {
-    Func("main", tint::Empty, ty.f32(),
-         Vector{
-             Return(1_f),
-         },
-         Vector{
-             Stage(ast::PipelineStage::kFragment),
-         },
-         Vector{
-             Location(Source{{12, 34}}, 2_a),
-             Location(Source{{56, 78}}, 3_a),
-         });
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              R"(56:78 error: duplicate location attribute
-12:34 note: first attribute declared here)");
-}
-
 TEST_F(EntryPointParameterAttributeTest, DuplicateInternalAttribute) {
     auto* s = Param("s", ty.sampler(core::type::SamplerKind::kSampler),
                     Vector{
@@ -625,25 +1194,6 @@
 }
 
 using EntryPointReturnTypeAttributeTest = ResolverTest;
-TEST_F(EntryPointReturnTypeAttributeTest, DuplicateAttribute) {
-    Func("main", tint::Empty, ty.f32(),
-         Vector{
-             Return(1_f),
-         },
-         Vector{
-             Stage(ast::PipelineStage::kFragment),
-         },
-         Vector{
-             Location(Source{{12, 34}}, 2_a),
-             Location(Source{{56, 78}}, 3_a),
-         });
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              R"(56:78 error: duplicate location attribute
-12:34 note: first attribute declared here)");
-}
-
 TEST_F(EntryPointReturnTypeAttributeTest, DuplicateInternalAttribute) {
     Func("f", tint::Empty, ty.i32(), Vector{Return(1_i)},
          Vector{
@@ -662,114 +1212,176 @@
 using StructAttributeTest = TestWithParams;
 using SpirvBlockAttribute = ast::transform::AddBlockAttribute::BlockAttribute;
 TEST_P(StructAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    Structure("mystruct", Vector{Member("a", ty.f32())},
-              createAttributes(Source{{12, 34}}, *this, params.kind));
+    Structure("S", Vector{Member("a", ty.f32())}, CreateAttributes());
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: " + name(params.kind) + " is not valid for struct declarations");
-    }
+    CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         StructAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    StructAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: @binding is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: @builtin is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: @diagnostic is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: @group is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: @id is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(1:2 error: @index is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: @interpolate is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: @invariant is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: @location is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: @stage is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is not valid for struct declarations)",
+        },
+        TestParams{
+            {AttributeKind::kBinding, AttributeKind::kGroup},
+            R"(1:2 error: @binding is not valid for struct declarations)",
+        }));
 
 using StructMemberAttributeTest = TestWithParams;
 TEST_P(StructMemberAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
-    Vector<const ast::StructMember*, 1> members;
-    if (params.kind == AttributeKind::kBuiltin) {
-        members.Push(
-            Member("a", ty.vec4<f32>(), createAttributes(Source{{12, 34}}, *this, params.kind)));
-    } else {
-        members.Push(Member("a", ty.f32(), createAttributes(Source{{12, 34}}, *this, params.kind)));
-    }
-    Structure("mystruct", members);
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: " + name(params.kind) + " is not valid for struct members");
-    }
+    EnableRequiredExtensions();
+    Structure("S", Vector{Member("a", ty.vec4<f32>(), CreateAttributes())});
+
+    CHECK();
 }
 INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                          StructMemberAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, true},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, true},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         // kIndex tested separately (requires @location)
-                                         // kInterpolate tested separately (requires @location)
-                                         // kInvariant tested separately (requires position builtin)
-                                         TestParams{AttributeKind::kLocation, true},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, true},
-                                         TestParams{AttributeKind::kSize, true},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
-TEST_F(StructMemberAttributeTest, DuplicateAttribute) {
-    Structure("mystruct", Vector{
-                              Member("a", ty.i32(),
-                                     Vector{
-                                         MemberAlign(Source{{12, 34}}, 4_i),
-                                         MemberAlign(Source{{56, 78}}, 8_i),
-                                     }),
-                          });
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              R"(56:78 error: duplicate align attribute
-12:34 note: first attribute declared here)");
-}
-TEST_F(StructMemberAttributeTest, InvariantAttributeWithPosition) {
-    Structure("mystruct", Vector{
-                              Member("a", ty.vec4<f32>(),
-                                     Vector{
-                                         Invariant(),
-                                         Builtin(core::BuiltinValue::kPosition),
-                                     }),
-                          });
-    EXPECT_TRUE(r()->Resolve()) << r()->error();
-}
-TEST_F(StructMemberAttributeTest, InvariantAttributeWithoutPosition) {
-    Structure("mystruct", Vector{
-                              Member("a", ty.vec4<f32>(),
-                                     Vector{
-                                         Invariant(Source{{12, 34}}),
-                                     }),
-                          });
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              "12:34 error: invariant attribute must only be applied to a "
-              "position builtin");
-}
+                         testing::Values(
+                             TestParams{
+                                 {AttributeKind::kAlign},
+                                 Pass,
+                             },
+                             TestParams{
+                                 {AttributeKind::kBinding},
+                                 R"(1:2 error: @binding is not valid for struct members)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kBuiltinPosition},
+                                 Pass,
+                             },
+                             TestParams{
+                                 {AttributeKind::kDiagnostic},
+                                 R"(1:2 error: @diagnostic is not valid for struct members)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kGroup},
+                                 R"(1:2 error: @group is not valid for struct members)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kId},
+                                 R"(1:2 error: @id is not valid for struct members)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kIndex},
+                                 R"(1:2 error: @index can only be used with @location(0))",
+                             },
+                             TestParams{
+                                 {AttributeKind::kInterpolate},
+                                 R"(1:2 error: @interpolate can only be used with @location)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kInterpolate, AttributeKind::kLocation},
+                                 Pass,
+                             },
+                             TestParams{
+                                 {AttributeKind::kInvariant},
+                                 R"(1:2 error: @invariant must be applied to a position builtin)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kInvariant, AttributeKind::kBuiltinPosition},
+                                 Pass,
+                             },
+                             TestParams{
+                                 {AttributeKind::kLocation},
+                                 Pass,
+                             },
+                             TestParams{
+                                 {AttributeKind::kMustUse},
+                                 R"(1:2 error: @must_use is not valid for struct members)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kOffset},
+                                 Pass,
+                             },
+                             TestParams{
+                                 {AttributeKind::kSize},
+                                 Pass,
+                             },
+                             TestParams{
+                                 {AttributeKind::kStageCompute},
+                                 R"(1:2 error: @stage is not valid for struct members)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kStride},
+                                 R"(1:2 error: @stride is not valid for struct members)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kWorkgroupSize},
+                                 R"(1:2 error: @workgroup_size is not valid for struct members)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kBinding, AttributeKind::kGroup},
+                                 R"(1:2 error: @binding is not valid for struct members)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kAlign, AttributeKind::kAlign},
+                                 R"(3:4 error: duplicate align attribute
+1:2 note: first attribute declared here)",
+                             }));
 
 TEST_F(StructMemberAttributeTest, Align_Attribute_Const) {
     GlobalConst("val", ty.i32(), Expr(1_i));
@@ -945,93 +1557,182 @@
 
 using ArrayAttributeTest = TestWithParams;
 TEST_P(ArrayAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    auto arr = ty.array(ty.f32(), createAttributes(Source{{12, 34}}, *this, params.kind));
-    Structure("mystruct", Vector{
-                              Member("a", arr),
-                          });
+    auto arr = ty.array(ty.f32(), CreateAttributes());
+    Structure("S", Vector{
+                       Member("a", arr),
+                   });
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: " + name(params.kind) + " is not valid for array types");
-    }
+    CHECK();
 }
 INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                          ArrayAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, true},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+                         testing::Values(
+                             TestParams{
+                                 {AttributeKind::kAlign},
+                                 R"(1:2 error: @align is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kBinding},
+                                 R"(1:2 error: @binding is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kBuiltinPosition},
+                                 R"(1:2 error: @builtin is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kDiagnostic},
+                                 R"(1:2 error: @diagnostic is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kGroup},
+                                 R"(1:2 error: @group is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kId},
+                                 R"(1:2 error: @id is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kIndex},
+                                 R"(1:2 error: @index is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kInterpolate},
+                                 R"(1:2 error: @interpolate is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kInvariant},
+                                 R"(1:2 error: @invariant is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kLocation},
+                                 R"(1:2 error: @location is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kMustUse},
+                                 R"(1:2 error: @must_use is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kOffset},
+                                 R"(1:2 error: @offset is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kSize},
+                                 R"(1:2 error: @size is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kStageCompute},
+                                 R"(1:2 error: @stage is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kStride},
+                                 Pass,
+                             },
+                             TestParams{
+                                 {AttributeKind::kWorkgroupSize},
+                                 R"(1:2 error: @workgroup_size is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kBinding, AttributeKind::kGroup},
+                                 R"(1:2 error: @binding is not valid for array types)",
+                             },
+                             TestParams{
+                                 {AttributeKind::kStride, AttributeKind::kStride},
+                                 R"(3:4 error: duplicate stride attribute
+1:2 note: first attribute declared here)",
+                             }));
 
 using VariableAttributeTest = TestWithParams;
 TEST_P(VariableAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    auto attrs = createAttributes(Source{{12, 34}}, *this, params.kind);
-    if (IsBindingAttribute(params.kind)) {
-        GlobalVar("a", ty.sampler(core::type::SamplerKind::kSampler), attrs);
+    if (GetParam().attributes.Any(IsBindingAttribute)) {
+        GlobalVar(Source{{9, 9}}, "a", ty.sampler(core::type::SamplerKind::kSampler),
+                  CreateAttributes());
     } else {
-        GlobalVar("a", ty.f32(), core::AddressSpace::kPrivate, attrs);
+        GlobalVar(Source{{9, 9}}, "a", ty.f32(), core::AddressSpace::kPrivate, CreateAttributes());
     }
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        if (!IsBindingAttribute(params.kind)) {
-            EXPECT_EQ(r()->error(),
-                      "12:34 error: " + name(params.kind) + " is not valid for module-scope 'var'");
-        }
-    }
+    CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         VariableAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, true}));
-
-TEST_F(VariableAttributeTest, DuplicateAttribute) {
-    GlobalVar("a", ty.sampler(core::type::SamplerKind::kSampler), Binding(Source{{12, 34}}, 2_a),
-              Group(2_a), Binding(Source{{56, 78}}, 3_a));
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              R"(56:78 error: duplicate binding attribute
-12:34 note: first attribute declared here)");
-}
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    VariableAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(9:9 error: resource variables require @group and @binding attributes)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: @builtin is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: @diagnostic is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(9:9 error: resource variables require @group and @binding attributes)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: @id is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(1:2 error: @index is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: @interpolate is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: @invariant is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: @location is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: @stage is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is not valid for module-scope 'var')",
+        },
+        TestParams{
+            {AttributeKind::kBinding, AttributeKind::kGroup},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kBinding, AttributeKind::kGroup, AttributeKind::kBinding},
+            R"(5:6 error: duplicate binding attribute
+1:2 note: first attribute declared here)",
+        }));
 
 TEST_F(VariableAttributeTest, LocalVar) {
     auto* v = Var("a", ty.f32(), Vector{Binding(Source{{12, 34}}, 2_a)});
@@ -1053,420 +1754,319 @@
 
 using ConstantAttributeTest = TestWithParams;
 TEST_P(ConstantAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    GlobalConst("a", ty.f32(), Expr(1.23_f),
-                createAttributes(Source{{12, 34}}, *this, params.kind));
+    GlobalConst("a", ty.f32(), Expr(1.23_f), CreateAttributes());
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: " + name(params.kind) + " is not valid for 'const' declaration");
-    }
+    CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         ConstantAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
-
-TEST_F(ConstantAttributeTest, InvalidAttribute) {
-    GlobalConst("a", ty.f32(), Expr(1.23_f),
-                Vector{
-                    Id(Source{{12, 34}}, 0_a),
-                });
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(), "12:34 error: @id is not valid for 'const' declaration");
-}
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    ConstantAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: @binding is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: @builtin is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: @diagnostic is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: @group is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            R"(1:2 error: @id is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(1:2 error: @index is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: @interpolate is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: @invariant is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: @location is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: @stage is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is not valid for 'const' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kBinding, AttributeKind::kGroup},
+            R"(1:2 error: @binding is not valid for 'const' declaration)",
+        }));
 
 using OverrideAttributeTest = TestWithParams;
 TEST_P(OverrideAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    Override("a", ty.f32(), Expr(1.23_f), createAttributes(Source{{12, 34}}, *this, params.kind));
+    Override("a", ty.f32(), Expr(1.23_f), CreateAttributes());
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: " + name(params.kind) + " is not valid for 'override' declaration");
-    }
+    CHECK();
 }
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
-                         OverrideAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, false},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kId, true},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
-
-TEST_F(OverrideAttributeTest, DuplicateAttribute) {
-    Override("a", ty.f32(), Expr(1.23_f),
-             Vector{
-                 Id(Source{{12, 34}}, 0_a),
-                 Id(Source{{56, 78}}, 1_a),
-             });
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              R"(56:78 error: duplicate id attribute
-12:34 note: first attribute declared here)");
-}
+INSTANTIATE_TEST_SUITE_P(
+    ResolverAttributeValidationTest,
+    OverrideAttributeTest,
+    testing::Values(
+        TestParams{
+            {AttributeKind::kAlign},
+            R"(1:2 error: @align is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kBinding},
+            R"(1:2 error: @binding is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kBuiltinPosition},
+            R"(1:2 error: @builtin is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kDiagnostic},
+            R"(1:2 error: @diagnostic is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kGroup},
+            R"(1:2 error: @group is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kIndex},
+            R"(1:2 error: @index is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kId},
+            Pass,
+        },
+        TestParams{
+            {AttributeKind::kInterpolate},
+            R"(1:2 error: @interpolate is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kInvariant},
+            R"(1:2 error: @invariant is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kLocation},
+            R"(1:2 error: @location is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kMustUse},
+            R"(1:2 error: @must_use is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kOffset},
+            R"(1:2 error: @offset is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kSize},
+            R"(1:2 error: @size is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kStageCompute},
+            R"(1:2 error: @stage is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kStride},
+            R"(1:2 error: @stride is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kWorkgroupSize},
+            R"(1:2 error: @workgroup_size is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kBinding, AttributeKind::kGroup},
+            R"(1:2 error: @binding is not valid for 'override' declaration)",
+        },
+        TestParams{
+            {AttributeKind::kId, AttributeKind::kId},
+            R"(3:4 error: duplicate id attribute
+1:2 note: first attribute declared here)",
+        }));
 
 using SwitchStatementAttributeTest = TestWithParams;
 TEST_P(SwitchStatementAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    WrapInFunction(Switch(Expr(0_a), Vector{DefaultCase()},
-                          createAttributes(Source{{12, 34}}, *this, params.kind)));
+    WrapInFunction(Switch(Expr(0_a), Vector{DefaultCase()}, CreateAttributes()));
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: " + name(params.kind) + " is not valid for switch statements");
-    }
+    CHECK();
 }
 INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                          SwitchStatementAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, true},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+                         testing::ValuesIn(OnlyDiagnosticValidFor("switch statements")));
 
 using SwitchBodyAttributeTest = TestWithParams;
 TEST_P(SwitchBodyAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    WrapInFunction(Switch(Expr(0_a), Vector{DefaultCase()}, tint::Empty,
-                          createAttributes(Source{{12, 34}}, *this, params.kind)));
+    WrapInFunction(Switch(Expr(0_a), Vector{DefaultCase()}, tint::Empty, CreateAttributes()));
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: " + name(params.kind) + " is not valid for switch body");
-    }
+    CHECK();
 }
 INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                          SwitchBodyAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, true},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+                         testing::ValuesIn(OnlyDiagnosticValidFor("switch body")));
 
 using IfStatementAttributeTest = TestWithParams;
 TEST_P(IfStatementAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    WrapInFunction(If(Expr(true), Block(), ElseStmt(),
-                      createAttributes(Source{{12, 34}}, *this, params.kind)));
+    WrapInFunction(If(Expr(true), Block(), ElseStmt(), CreateAttributes()));
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: " + name(params.kind) + " is not valid for if statements");
-    }
+    CHECK();
 }
 INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                          IfStatementAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, true},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+                         testing::ValuesIn(OnlyDiagnosticValidFor("if statements")));
 
 using ForStatementAttributeTest = TestWithParams;
 TEST_P(ForStatementAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    WrapInFunction(For(nullptr, Expr(false), nullptr, Block(),
-                       createAttributes(Source{{12, 34}}, *this, params.kind)));
+    WrapInFunction(For(nullptr, Expr(false), nullptr, Block(), CreateAttributes()));
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: " + name(params.kind) + " is not valid for for statements");
-    }
+    CHECK();
 }
 INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                          ForStatementAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, true},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+                         testing::ValuesIn(OnlyDiagnosticValidFor("for statements")));
 
 using LoopStatementAttributeTest = TestWithParams;
 TEST_P(LoopStatementAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    WrapInFunction(
-        Loop(Block(Return()), Block(), createAttributes(Source{{12, 34}}, *this, params.kind)));
+    WrapInFunction(Loop(Block(Return()), Block(), CreateAttributes()));
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: " + name(params.kind) + " is not valid for loop statements");
-    }
+    CHECK();
 }
 INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                          LoopStatementAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, true},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+                         testing::ValuesIn(OnlyDiagnosticValidFor("loop statements")));
 
 using WhileStatementAttributeTest = TestWithParams;
 TEST_P(WhileStatementAttributeTest, IsValid) {
-    auto& params = GetParam();
-    EnableExtensionIfNecessary(params.kind);
+    EnableRequiredExtensions();
 
-    WrapInFunction(
-        While(Expr(false), Block(), createAttributes(Source{{12, 34}}, *this, params.kind)));
+    WrapInFunction(While(Expr(false), Block(), CreateAttributes()));
 
-    if (params.should_pass) {
-        EXPECT_TRUE(r()->Resolve()) << r()->error();
-    } else {
-        EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: " + name(params.kind) + " is not valid for while statements");
-    }
+    CHECK();
 }
 INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                          WhileStatementAttributeTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, true},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
+                         testing::ValuesIn(OnlyDiagnosticValidFor("while statements")));
 
-namespace BlockStatementTests {
-class BlockStatementTest : public TestWithParams {
-  protected:
-    void Check() {
-        if (GetParam().should_pass) {
-            EXPECT_TRUE(r()->Resolve()) << r()->error();
-        } else {
-            EXPECT_FALSE(r()->Resolve());
-            EXPECT_EQ(r()->error(),
-                      "error: " + name(GetParam().kind) + " is not valid for block statements");
-        }
-    }
-
-  public:
-    BlockStatementTest() { EnableExtensionIfNecessary(GetParam().kind); }
-};
+using BlockStatementTest = TestWithParams;
 TEST_P(BlockStatementTest, CompoundStatement) {
     Func("foo", tint::Empty, ty.void_(),
          Vector{
-             Block(Vector{Return()}, createAttributes({}, *this, GetParam().kind)),
+             Block(Vector{Return()}, CreateAttributes()),
          });
-    Check();
+
+    CHECK();
 }
 TEST_P(BlockStatementTest, FunctionBody) {
-    Func("foo", tint::Empty, ty.void_(),
-         Block(Vector{Return()}, createAttributes({}, *this, GetParam().kind)));
-    Check();
+    Func("foo", tint::Empty, ty.void_(), Block(Vector{Return()}, CreateAttributes()));
+
+    CHECK();
 }
 TEST_P(BlockStatementTest, IfStatementBody) {
     Func("foo", tint::Empty, ty.void_(),
          Vector{
-             If(Expr(true), Block(Vector{Return()}, createAttributes({}, *this, GetParam().kind))),
+             If(Expr(true), Block(Vector{Return()}, CreateAttributes())),
          });
-    Check();
+
+    CHECK();
 }
 TEST_P(BlockStatementTest, ElseStatementBody) {
     Func("foo", tint::Empty, ty.void_(),
          Vector{
              If(Expr(true), Block(Vector{Return()}),
-                Else(Block(Vector{Return()}, createAttributes({}, *this, GetParam().kind)))),
+                Else(Block(Vector{Return()}, CreateAttributes()))),
          });
-    Check();
+
+    CHECK();
 }
 TEST_P(BlockStatementTest, ForStatementBody) {
     Func("foo", tint::Empty, ty.void_(),
          Vector{
-             For(nullptr, Expr(true), nullptr,
-                 Block(Vector{Break()}, createAttributes({}, *this, GetParam().kind))),
+             For(nullptr, Expr(true), nullptr, Block(Vector{Break()}, CreateAttributes())),
          });
-    Check();
+
+    CHECK();
 }
 TEST_P(BlockStatementTest, LoopStatementBody) {
     Func("foo", tint::Empty, ty.void_(),
          Vector{
-             Loop(Block(Vector{Break()}, createAttributes({}, *this, GetParam().kind))),
+             Loop(Block(Vector{Break()}, CreateAttributes())),
          });
-    Check();
+
+    CHECK();
 }
 TEST_P(BlockStatementTest, WhileStatementBody) {
-    Func(
-        "foo", tint::Empty, ty.void_(),
-        Vector{
-            While(Expr(true), Block(Vector{Break()}, createAttributes({}, *this, GetParam().kind))),
-        });
-    Check();
+    Func("foo", tint::Empty, ty.void_(),
+         Vector{
+             While(Expr(true), Block(Vector{Break()}, CreateAttributes())),
+         });
+
+    CHECK();
 }
 TEST_P(BlockStatementTest, CaseStatementBody) {
     Func("foo", tint::Empty, ty.void_(),
          Vector{
-             Switch(1_a,
-                    Case(CaseSelector(1_a),
-                         Block(Vector{Break()}, createAttributes({}, *this, GetParam().kind))),
+             Switch(1_a, Case(CaseSelector(1_a), Block(Vector{Break()}, CreateAttributes())),
                     DefaultCase(Block({}))),
          });
-    Check();
+
+    CHECK();
 }
 TEST_P(BlockStatementTest, DefaultStatementBody) {
     Func("foo", tint::Empty, ty.void_(),
          Vector{
-             Switch(
-                 1_a, Case(CaseSelector(1_a), Block()),
-                 DefaultCase(Block(Vector{Break()}, createAttributes({}, *this, GetParam().kind)))),
+             Switch(1_a, Case(CaseSelector(1_a), Block()),
+                    DefaultCase(Block(Vector{Break()}, CreateAttributes()))),
          });
-    Check();
+
+    CHECK();
 }
 INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                          BlockStatementTest,
-                         testing::Values(TestParams{AttributeKind::kAlign, false},
-                                         TestParams{AttributeKind::kBinding, false},
-                                         TestParams{AttributeKind::kBuiltin, false},
-                                         TestParams{AttributeKind::kDiagnostic, true},
-                                         TestParams{AttributeKind::kGroup, false},
-                                         TestParams{AttributeKind::kId, false},
-                                         TestParams{AttributeKind::kIndex, false},
-                                         TestParams{AttributeKind::kInterpolate, false},
-                                         TestParams{AttributeKind::kInvariant, false},
-                                         TestParams{AttributeKind::kLocation, false},
-                                         TestParams{AttributeKind::kMustUse, false},
-                                         TestParams{AttributeKind::kOffset, false},
-                                         TestParams{AttributeKind::kSize, false},
-                                         TestParams{AttributeKind::kStage, false},
-                                         TestParams{AttributeKind::kStride, false},
-                                         TestParams{AttributeKind::kWorkgroup, false},
-                                         TestParams{AttributeKind::kBindingAndGroup, false}));
-
-}  // namespace BlockStatementTests
+                         testing::ValuesIn(OnlyDiagnosticValidFor("block statements")));
 
 }  // namespace
 }  // namespace AttributeTests
@@ -1509,9 +2109,9 @@
     } else {
         EXPECT_FALSE(r()->Resolve());
         EXPECT_EQ(r()->error(),
-                  "12:34 error: arrays decorated with the stride attribute must "
-                  "have a stride that is at least the size of the element type, "
-                  "and be a multiple of the element type's alignment value");
+                  "12:34 error: arrays decorated with the stride attribute must have a stride that "
+                  "is at least the size of the element type, and be a multiple of the element "
+                  "type's alignment value");
     }
 }
 
@@ -1720,96 +2320,10 @@
 }  // namespace
 }  // namespace ResourceTests
 
-namespace InvariantAttributeTests {
-namespace {
-using InvariantAttributeTests = ResolverTest;
-TEST_F(InvariantAttributeTests, InvariantWithPosition) {
-    auto* param = Param("p", ty.vec4<f32>(),
-                        Vector{
-                            Invariant(Source{{12, 34}}),
-                            Builtin(Source{{56, 78}}, core::BuiltinValue::kPosition),
-                        });
-    Func("main", Vector{param}, ty.vec4<f32>(),
-         Vector{
-             Return(Call<vec4<f32>>()),
-         },
-         Vector{
-             Stage(ast::PipelineStage::kFragment),
-         },
-         Vector{
-             Location(0_a),
-         });
-    EXPECT_TRUE(r()->Resolve()) << r()->error();
-}
-
-TEST_F(InvariantAttributeTests, InvariantWithoutPosition) {
-    auto* param = Param("p", ty.vec4<f32>(),
-                        Vector{
-                            Invariant(Source{{12, 34}}),
-                            Location(0_a),
-                        });
-    Func("main", Vector{param}, ty.vec4<f32>(),
-         Vector{
-             Return(Call<vec4<f32>>()),
-         },
-         Vector{
-             Stage(ast::PipelineStage::kFragment),
-         },
-         Vector{
-             Location(0_a),
-         });
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              "12:34 error: invariant attribute must only be applied to a "
-              "position builtin");
-}
-}  // namespace
-}  // namespace InvariantAttributeTests
-
-namespace MustUseAttributeTests {
-namespace {
-
-using MustUseAttributeTests = ResolverTest;
-TEST_F(MustUseAttributeTests, MustUse) {
-    Func("main", tint::Empty, ty.vec4<f32>(),
-         Vector{
-             Return(Call<vec4<f32>>()),
-         },
-         Vector{
-             MustUse(Source{{12, 34}}),
-         });
-    EXPECT_TRUE(r()->Resolve()) << r()->error();
-}
-
-}  // namespace
-}  // namespace MustUseAttributeTests
-
 namespace WorkgroupAttributeTests {
 namespace {
 
 using WorkgroupAttribute = ResolverTest;
-TEST_F(WorkgroupAttribute, ComputeShaderPass) {
-    Func("main", tint::Empty, ty.void_(), tint::Empty,
-         Vector{
-             Stage(ast::PipelineStage::kCompute),
-             create<ast::WorkgroupAttribute>(Source{{12, 34}}, Expr(1_i)),
-         });
-
-    EXPECT_TRUE(r()->Resolve()) << r()->error();
-}
-
-TEST_F(WorkgroupAttribute, Missing) {
-    Func(Source{{12, 34}}, "main", tint::Empty, ty.void_(), tint::Empty,
-         Vector{
-             Stage(ast::PipelineStage::kCompute),
-         });
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              "12:34 error: a compute shader must include 'workgroup_size' in its "
-              "attributes");
-}
-
 TEST_F(WorkgroupAttribute, NotAnEntryPoint) {
     Func("main", tint::Empty, ty.void_(), tint::Empty,
          Vector{
@@ -1817,7 +2331,7 @@
          });
 
     EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(), "12:34 error: @workgroup_size is only valid for compute stages");
+    EXPECT_EQ(r()->error(), R"(12:34 error: @workgroup_size is only valid for compute stages)");
 }
 
 TEST_F(WorkgroupAttribute, NotAComputeShader) {
@@ -1828,7 +2342,7 @@
          });
 
     EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(), "12:34 error: @workgroup_size is only valid for compute stages");
+    EXPECT_EQ(r()->error(), R"(12:34 error: @workgroup_size is only valid for compute stages)");
 }
 
 TEST_F(WorkgroupAttribute, DuplicateAttribute) {
@@ -1864,7 +2378,6 @@
 using InterpolateParameterTest = TestWithParams;
 TEST_P(InterpolateParameterTest, All) {
     auto& params = GetParam();
-
     Func("main",
          Vector{
              Param("a", ty.f32(),
@@ -1882,15 +2395,14 @@
         EXPECT_TRUE(r()->Resolve()) << r()->error();
     } else {
         EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: flat interpolation attribute must not have a "
-                  "sampling parameter");
+        EXPECT_EQ(
+            r()->error(),
+            R"(12:34 error: flat interpolation attribute must not have a sampling parameter)");
     }
 }
 
 TEST_P(InterpolateParameterTest, IntegerScalar) {
     auto& params = GetParam();
-
     Func("main",
          Vector{
              Param("a", ty.i32(),
@@ -1906,22 +2418,21 @@
 
     if (params.type != core::InterpolationType::kFlat) {
         EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: interpolation type must be 'flat' for integral "
-                  "user-defined IO types");
+        EXPECT_EQ(
+            r()->error(),
+            R"(12:34 error: interpolation type must be 'flat' for integral user-defined IO types)");
     } else if (params.should_pass) {
         EXPECT_TRUE(r()->Resolve()) << r()->error();
     } else {
         EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: flat interpolation attribute must not have a "
-                  "sampling parameter");
+        EXPECT_EQ(
+            r()->error(),
+            R"(12:34 error: flat interpolation attribute must not have a sampling parameter)");
     }
 }
 
 TEST_P(InterpolateParameterTest, IntegerVector) {
     auto& params = GetParam();
-
     Func("main",
          Vector{
              Param("a", ty.vec4<u32>(),
@@ -1937,16 +2448,16 @@
 
     if (params.type != core::InterpolationType::kFlat) {
         EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: interpolation type must be 'flat' for integral "
-                  "user-defined IO types");
+        EXPECT_EQ(
+            r()->error(),
+            R"(12:34 error: interpolation type must be 'flat' for integral user-defined IO types)");
     } else if (params.should_pass) {
         EXPECT_TRUE(r()->Resolve()) << r()->error();
     } else {
         EXPECT_FALSE(r()->Resolve());
-        EXPECT_EQ(r()->error(),
-                  "12:34 error: flat interpolation attribute must not have a "
-                  "sampling parameter");
+        EXPECT_EQ(
+            r()->error(),
+            R"(12:34 error: flat interpolation attribute must not have a sampling parameter)");
     }
 }
 
@@ -2003,55 +2514,6 @@
 note: while analyzing entry point 'main')");
 }
 
-TEST_F(InterpolateTest, MissingLocationAttribute_Parameter) {
-    Func("main",
-         Vector{
-             Param("a", ty.vec4<f32>(),
-                   Vector{
-                       Builtin(core::BuiltinValue::kPosition),
-                       Interpolate(Source{{12, 34}}, core::InterpolationType::kFlat),
-                   }),
-         },
-         ty.void_(), tint::Empty,
-         Vector{
-             Stage(ast::PipelineStage::kFragment),
-         });
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              R"(12:34 error: interpolate attribute must only be used with @location)");
-}
-
-TEST_F(InterpolateTest, MissingLocationAttribute_ReturnType) {
-    Func("main", tint::Empty, ty.vec4<f32>(),
-         Vector{
-             Return(Call<vec4<f32>>()),
-         },
-         Vector{
-             Stage(ast::PipelineStage::kVertex),
-         },
-         Vector{
-             Builtin(core::BuiltinValue::kPosition),
-             Interpolate(Source{{12, 34}}, core::InterpolationType::kFlat),
-         });
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              R"(12:34 error: interpolate attribute must only be used with @location)");
-}
-
-TEST_F(InterpolateTest, MissingLocationAttribute_Struct) {
-    Structure("S",
-              Vector{
-                  Member("a", ty.f32(),
-                         Vector{Interpolate(Source{{12, 34}}, core::InterpolationType::kFlat)}),
-              });
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              R"(12:34 error: interpolate attribute must only be used with @location)");
-}
-
 using GroupAndBindingTest = ResolverTest;
 
 TEST_F(GroupAndBindingTest, Const_I32) {
@@ -2287,21 +2749,6 @@
 }  // namespace
 }  // namespace InterpolateTests
 
-namespace MustUseTests {
-namespace {
-
-using MustUseAttributeTest = ResolverTest;
-TEST_F(MustUseAttributeTest, UsedOnFnWithNoReturnValue) {
-    Func("fn_must_use", tint::Empty, ty.void_(), tint::Empty, Vector{MustUse(Source{{12, 34}})});
-
-    EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              R"(12:34 error: @must_use can only be applied to functions that return a value)");
-}
-
-}  // namespace
-}  // namespace MustUseTests
-
 namespace InternalAttributeDeps {
 namespace {
 
diff --git a/src/tint/lang/wgsl/resolver/builtins_validation_test.cc b/src/tint/lang/wgsl/resolver/builtins_validation_test.cc
index 1da3470..616161d 100644
--- a/src/tint/lang/wgsl/resolver/builtins_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/builtins_validation_test.cc
@@ -147,7 +147,7 @@
     } else {
         StringStream err;
         err << "12:34 error: @builtin(" << params.builtin << ")";
-        err << " cannot be used in input of " << params.stage << " pipeline stage";
+        err << " cannot be used for " << params.stage << " shader input";
         EXPECT_FALSE(r()->Resolve());
         EXPECT_EQ(r()->error(), err.str());
     }
@@ -179,9 +179,8 @@
              Location(0_a),
          });
     EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(
-        r()->error(),
-        "12:34 error: @builtin(frag_depth) cannot be used in input of fragment pipeline stage");
+    EXPECT_EQ(r()->error(),
+              "12:34 error: @builtin(frag_depth) cannot be used for fragment shader input");
 }
 
 TEST_F(ResolverBuiltinsValidationTest, FragDepthIsInputStruct_Fail) {
@@ -215,9 +214,8 @@
          });
     EXPECT_FALSE(r()->Resolve());
     EXPECT_EQ(r()->error(),
-              "12:34 error: @builtin(frag_depth) cannot be used in input of "
-              "fragment pipeline stage\n"
-              "note: while analyzing entry point 'fragShader'");
+              R"(12:34 error: @builtin(frag_depth) cannot be used for fragment shader input
+note: while analyzing entry point 'fragShader')");
 }
 
 TEST_F(ResolverBuiltinsValidationTest, StructBuiltinInsideEntryPoint_Ignored) {
diff --git a/src/tint/lang/wgsl/resolver/dual_source_blending_extension_test.cc b/src/tint/lang/wgsl/resolver/dual_source_blending_extension_test.cc
index cb4489f..9ea0b05 100644
--- a/src/tint/lang/wgsl/resolver/dual_source_blending_extension_test.cc
+++ b/src/tint/lang/wgsl/resolver/dual_source_blending_extension_test.cc
@@ -47,9 +47,9 @@
               });
 
     EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(),
-              "12:34 error: use of '@index' attribute requires enabling extension "
-              "'chromium_internal_dual_source_blending'");
+    EXPECT_EQ(
+        r()->error(),
+        R"(12:34 error: use of @index requires enabling extension 'chromium_internal_dual_source_blending')");
 }
 
 class DualSourceBlendingExtensionTests : public ResolverTest {
@@ -107,7 +107,7 @@
     Structure("Output", Vector{
                             Member("a", ty.vec4<f32>(), Vector{Location(0_a), Index(0_a)}),
                             Member(Source{{12, 34}}, "b", ty.vec4<f32>(),
-                                   Vector{Location(0_a), Index(Source{{12, 34}}, 0_a)}),
+                                   Vector{Location(Source{{12, 34}}, 0_a), Index(0_a)}),
                         });
 
     EXPECT_FALSE(r()->Resolve());
@@ -115,14 +115,30 @@
 }
 
 // Using the index attribute without a location attribute should fail.
-TEST_F(DualSourceBlendingExtensionTests, IndexWithMissingLocationAttribute) {
+TEST_F(DualSourceBlendingExtensionTests, IndexWithMissingLocationAttribute_Struct) {
     Structure("Output", Vector{
                             Member(Source{{12, 34}}, "a", ty.vec4<f32>(),
                                    Vector{Index(Source{{12, 34}}, 1_a)}),
                         });
 
     EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(), "12:34 error: index attribute must only be used with @location");
+    EXPECT_EQ(r()->error(), "12:34 error: @index can only be used with @location(0)");
+}
+
+// Using the index attribute without a location attribute should fail.
+TEST_F(DualSourceBlendingExtensionTests, IndexWithMissingLocationAttribute_ReturnValue) {
+    Func("F", Empty, ty.vec4<f32>(),
+         Vector{
+             Return(Call<vec4<f32>>()),
+         },
+         Vector{Stage(ast::PipelineStage::kFragment)},
+         Vector{
+             Index(Source{{12, 34}}, 1_a),
+             Builtin(core::BuiltinValue::kPointSize),
+         });
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "12:34 error: @index can only be used with @location(0)");
 }
 
 // Using an index attribute on a struct member should pass.
@@ -148,14 +164,30 @@
 }
 
 // Using the index attribute with a non-zero location should fail.
-TEST_F(DualSourceBlendingExtensionTests, IndexWithNonZeroLocation) {
+TEST_F(DualSourceBlendingExtensionTests, IndexWithNonZeroLocation_Struct) {
     Structure("Output",
               Vector{
                   Member("a", ty.vec4<f32>(), Vector{Location(1_a), Index(Source{{12, 34}}, 0_a)}),
               });
 
     EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(), "12:34 error: index attribute must only be used with @location(0)");
+    EXPECT_EQ(r()->error(), "12:34 error: @index can only be used with @location(0)");
+}
+
+// Using the index attribute with a non-zero location should fail.
+TEST_F(DualSourceBlendingExtensionTests, IndexWithNonZeroLocation_ReturnValue) {
+    Func("F", Empty, ty.vec4<f32>(),
+         Vector{
+             Return(Call<vec4<f32>>()),
+         },
+         Vector{Stage(ast::PipelineStage::kFragment)},
+         Vector{
+             Location(1_a),
+             Index(Source{{12, 34}}, 1_a),
+         });
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "12:34 error: @index can only be used with @location(0)");
 }
 
 class DualSourceBlendingExtensionTestWithParams : public ResolverTestWithParam<int> {
@@ -166,20 +198,40 @@
 };
 
 // Rendering to multiple render targets while using dual source blending should fail.
-TEST_P(DualSourceBlendingExtensionTestWithParams, MultipleRenderTargetsNotAllowed) {
-    Structure("Output",
+TEST_P(DualSourceBlendingExtensionTestWithParams,
+       MultipleRenderTargetsNotAllowed_IndexThenNonZeroLocation) {
+    Structure("S",
               Vector{
                   Member("a", ty.vec4<f32>(), Vector{Location(0_a), Index(0_a)}),
-                  Member("b", ty.vec4<f32>(), Vector{Location(0_a), Index(1_a)}),
-                  Member("c", ty.vec4<f32>(), Vector{Location(Source{{12, 34}}, AInt(GetParam()))}),
+                  Member("b", ty.vec4<f32>(), Vector{Location(0_a), Index(Source{{1, 2}}, 1_a)}),
+                  Member("c", ty.vec4<f32>(), Vector{Location(Source{{3, 4}}, AInt(GetParam()))}),
               });
+    Func("F", Empty, ty("S"), Vector{Return(Call("S"))},
+         Vector{Stage(ast::PipelineStage::kFragment)});
 
     EXPECT_FALSE(r()->Resolve());
-    StringStream err;
-    err << "12:34 error: Multiple render targets are not allowed when using dual source blending. "
-           "The output @location("
-        << GetParam() << ") is not allowed as a render target.";
-    EXPECT_EQ(r()->error(), err.str());
+    EXPECT_EQ(r()->error(),
+              R"(1:2 error: pipeline cannot use both non-zero @index and non-zero @location
+3:4 note: non-zero @location declared here
+note: while analyzing entry point 'F')");
+}
+
+TEST_P(DualSourceBlendingExtensionTestWithParams,
+       MultipleRenderTargetsNotAllowed_NonZeroLocationThenIndex) {
+    Structure("S",
+              Vector{
+                  Member("a", ty.vec4<f32>(), Vector{Location(Source{{1, 2}}, AInt(GetParam()))}),
+                  Member("b", ty.vec4<f32>(), Vector{Location(0_a), Index(0_a)}),
+                  Member("c", ty.vec4<f32>(), Vector{Location(0_a), Index(Source{{3, 4}}, 1_a)}),
+              });
+    Func(Source{{5, 6}}, "F", Empty, ty("S"), Vector{Return(Call("S"))},
+         Vector{Stage(ast::PipelineStage::kFragment)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              R"(3:4 error: pipeline cannot use both non-zero @index and non-zero @location
+1:2 note: non-zero @location declared here
+5:6 note: while analyzing entry point 'F')");
 }
 
 INSTANTIATE_TEST_SUITE_P(DualSourceBlendingExtensionTests,
diff --git a/src/tint/lang/wgsl/resolver/entry_point_validation_test.cc b/src/tint/lang/wgsl/resolver/entry_point_validation_test.cc
index 034e3b0..444f9e5 100644
--- a/src/tint/lang/wgsl/resolver/entry_point_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/entry_point_validation_test.cc
@@ -1072,7 +1072,7 @@
          });
 
     EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(), R"(12:34 error: @location is not valid for compute shader output)");
+    EXPECT_EQ(r()->error(), R"(12:34 error: @location cannot be used by compute shaders)");
 }
 
 TEST_F(LocationAttributeTests, ComputeShaderLocation_Output) {
@@ -1087,7 +1087,7 @@
          });
 
     EXPECT_FALSE(r()->Resolve());
-    EXPECT_EQ(r()->error(), R"(12:34 error: @location is not valid for compute shader inputs)");
+    EXPECT_EQ(r()->error(), R"(12:34 error: @location cannot be used by compute shaders)");
 }
 
 TEST_F(LocationAttributeTests, ComputeShaderLocationStructMember_Output) {
@@ -1107,7 +1107,7 @@
 
     EXPECT_FALSE(r()->Resolve());
     EXPECT_EQ(r()->error(),
-              "12:34 error: @location is not valid for compute shader output\n"
+              "12:34 error: @location cannot be used by compute shaders\n"
               "56:78 note: while analyzing entry point 'main'");
 }
 
@@ -1126,7 +1126,7 @@
 
     EXPECT_FALSE(r()->Resolve());
     EXPECT_EQ(r()->error(),
-              "12:34 error: @location is not valid for compute shader inputs\n"
+              "12:34 error: @location cannot be used by compute shaders\n"
               "56:78 note: while analyzing entry point 'main'");
 }
 
diff --git a/src/tint/lang/wgsl/resolver/override_test.cc b/src/tint/lang/wgsl/resolver/override_test.cc
index 6bd7223..3077e79 100644
--- a/src/tint/lang/wgsl/resolver/override_test.cc
+++ b/src/tint/lang/wgsl/resolver/override_test.cc
@@ -46,7 +46,7 @@
         ASSERT_NE(sem, nullptr);
         EXPECT_EQ(sem->Declaration(), var);
         EXPECT_TRUE(sem->Declaration()->Is<ast::Override>());
-        EXPECT_EQ(sem->OverrideId().value, id);
+        EXPECT_EQ(sem->Attributes().override_id->value, id);
         EXPECT_FALSE(sem->ConstantValue());
     }
 };
diff --git a/src/tint/lang/wgsl/resolver/resolver.cc b/src/tint/lang/wgsl/resolver/resolver.cc
index c3d2f53..d5a0e93 100644
--- a/src/tint/lang/wgsl/resolver/resolver.cc
+++ b/src/tint/lang/wgsl/resolver/resolver.cc
@@ -389,7 +389,7 @@
                 }
 
                 auto o = OverrideId{static_cast<decltype(OverrideId::value)>(value)};
-                sem->SetOverrideId(o);
+                sem->Attributes().override_id = o;
 
                 // Track the constant IDs that are specified in the shader.
                 override_ids_.Add(o, sem);
@@ -626,7 +626,7 @@
                     if (!value) {
                         return kErrored;
                     }
-                    global->SetLocation(value.Get());
+                    global->Attributes().location = value.Get();
                     return kSuccess;
                 },
                 [&](const ast::IndexAttribute* attr) {
@@ -637,7 +637,7 @@
                     if (!value) {
                         return kErrored;
                     }
-                    global->SetIndex(value.Get());
+                    global->Attributes().index = value.Get();
                     return kSuccess;
                 },
                 [&](const ast::BuiltinAttribute* attr) {
@@ -675,7 +675,7 @@
         }
 
         if (group && binding) {
-            global->SetBindingPoint(BindingPoint{group.value(), binding.value()});
+            global->Attributes().binding_point = BindingPoint{group.value(), binding.value()};
         }
 
     } else {
@@ -720,7 +720,7 @@
                     if (TINT_UNLIKELY(!value)) {
                         return false;
                     }
-                    sem->SetLocation(value.Get());
+                    sem->Attributes().location = value.Get();
                     return true;
                 },
                 [&](const ast::BuiltinAttribute* attr) -> bool { return BuiltinAttribute(attr); },
@@ -766,7 +766,7 @@
             }
         }
         if (group && binding) {
-            sem->SetBindingPoint(BindingPoint{group.value(), binding.value()});
+            sem->Attributes().binding_point = BindingPoint{group.value(), binding.value()};
         }
     } else {
         for (auto* attribute : param->attributes) {
@@ -862,8 +862,8 @@
         auto* sem = sem_.Get(override);
 
         OverrideId id;
-        if (ast::HasAttribute<ast::IdAttribute>(override->attributes)) {
-            id = sem->OverrideId();
+        if (auto sem_id = sem->Attributes().override_id) {
+            id = *sem_id;
         } else {
             // No ID was specified, so allocate the next available ID.
             while (!ids_exhausted && override_ids_.Contains(next_id)) {
@@ -879,7 +879,7 @@
             increment_next_id();
         }
 
-        const_cast<sem::GlobalVariable*>(sem)->SetOverrideId(id);
+        const_cast<sem::GlobalVariable*>(sem)->Attributes().override_id = id;
     }
     return true;
 }
diff --git a/src/tint/lang/wgsl/resolver/resolver_test.cc b/src/tint/lang/wgsl/resolver/resolver_test.cc
index 9b3c171..f78bdc3 100644
--- a/src/tint/lang/wgsl/resolver/resolver_test.cc
+++ b/src/tint/lang/wgsl/resolver/resolver_test.cc
@@ -896,9 +896,9 @@
     auto* func_sem = Sem().Get(func);
     ASSERT_NE(func_sem, nullptr);
     EXPECT_EQ(func_sem->Parameters().Length(), 3u);
-    EXPECT_EQ(3u, func_sem->Parameters()[0]->Location());
-    EXPECT_FALSE(func_sem->Parameters()[1]->Location().has_value());
-    EXPECT_EQ(1u, func_sem->Parameters()[2]->Location());
+    EXPECT_EQ(3u, func_sem->Parameters()[0]->Attributes().location);
+    EXPECT_FALSE(func_sem->Parameters()[1]->Attributes().location.has_value());
+    EXPECT_EQ(1u, func_sem->Parameters()[2]->Attributes().location);
 }
 
 TEST_F(ResolverTest, Function_GlobalVariable_Location) {
@@ -910,7 +910,7 @@
 
     auto* sem = Sem().Get<sem::GlobalVariable>(var);
     ASSERT_NE(sem, nullptr);
-    EXPECT_EQ(3u, sem->Location());
+    EXPECT_EQ(3u, sem->Attributes().location);
 }
 
 TEST_F(ResolverTest, Function_RegisterInputOutputVariables) {
@@ -1931,8 +1931,10 @@
 
     EXPECT_TRUE(r()->Resolve()) << r()->error();
 
-    EXPECT_EQ(Sem().Get<sem::GlobalVariable>(s1)->BindingPoint(), (BindingPoint{1u, 2u}));
-    EXPECT_EQ(Sem().Get<sem::GlobalVariable>(s2)->BindingPoint(), (BindingPoint{3u, 4u}));
+    EXPECT_EQ(Sem().Get<sem::GlobalVariable>(s1)->Attributes().binding_point,
+              (BindingPoint{1u, 2u}));
+    EXPECT_EQ(Sem().Get<sem::GlobalVariable>(s2)->Attributes().binding_point,
+              (BindingPoint{3u, 4u}));
 }
 
 TEST_F(ResolverTest, Function_EntryPoints_StageAttribute) {
diff --git a/src/tint/lang/wgsl/resolver/validator.cc b/src/tint/lang/wgsl/resolver/validator.cc
index c05eca2..67ba42d 100644
--- a/src/tint/lang/wgsl/resolver/validator.cc
+++ b/src/tint/lang/wgsl/resolver/validator.cc
@@ -137,10 +137,7 @@
 
 // Helper to stringify a pipeline IO attribute.
 std::string AttrToStr(const ast::Attribute* attr) {
-    return Switch(
-        attr,  //
-        [&](const ast::BuiltinAttribute*) { return "@builtin"; },
-        [&](const ast::LocationAttribute*) { return "@location"; });
+    return "@" + attr->Name();
 }
 
 template <typename CALLBACK>
@@ -770,17 +767,14 @@
         return false;
     }
 
-    for (auto* attr : decl->attributes) {
-        if (attr->Is<ast::IdAttribute>()) {
-            auto id = v->OverrideId();
-            if (auto var = override_ids.Find(id); var && *var != v) {
-                AddError("@id values must be unique", attr->source);
-                AddNote(
-                    "a override with an ID of " + std::to_string(id.value) +
+    if (auto id = v->Attributes().override_id) {
+        if (auto var = override_ids.Find(*id); var && *var != v) {
+            auto* attr = ast::GetAttribute<ast::IdAttribute>(v->Declaration()->attributes);
+            AddError("@id values must be unique", attr->source);
+            AddNote("a override with an ID of " + std::to_string(id->value) +
                         " was previously declared here:",
                     ast::GetAttribute<ast::IdAttribute>((*var)->Declaration()->attributes)->source);
-                return false;
-            }
+            return false;
         }
     }
 
@@ -990,8 +984,8 @@
 
     if (is_stage_mismatch) {
         StringStream err;
-        err << "@builtin(" << builtin << ") cannot be used in "
-            << (is_input ? "input of " : "output of ") << stage_name.str() << " pipeline stage";
+        err << "@builtin(" << builtin << ") cannot be used for " << stage_name.str() << " shader "
+            << (is_input ? "input" : "output");
         AddError(err.str(), attr->source);
         return false;
     }
@@ -1000,7 +994,13 @@
 }
 
 bool Validator::InterpolateAttribute(const ast::InterpolateAttribute* attr,
-                                     const core::type::Type* storage_ty) const {
+                                     const core::type::Type* storage_ty,
+                                     const ast::PipelineStage stage) const {
+    if (stage == ast::PipelineStage::kCompute) {
+        AddError(AttrToStr(attr) + " cannot be used by compute shaders", attr->source);
+        return false;
+    }
+
     auto* type = storage_ty->UnwrapRef();
 
     auto i_type = sem_.AsInterpolationType(sem_.Get(attr->type));
@@ -1022,6 +1022,15 @@
     return true;
 }
 
+bool Validator::InvariantAttribute(const ast::InvariantAttribute* attr,
+                                   const ast::PipelineStage stage) const {
+    if (stage == ast::PipelineStage::kCompute) {
+        AddError(AttrToStr(attr) + " cannot be used by compute shaders", attr->source);
+        return false;
+    }
+    return true;
+}
+
 bool Validator::Function(const sem::Function* func, ast::PipelineStage stage) const {
     auto* decl = func->Declaration();
 
@@ -1104,7 +1113,9 @@
     // TODO(jrprice): This state could be stored in sem::Function instead, and then passed to
     // sem::Function since it would be useful there too.
     Hashset<core::BuiltinValue, 4> builtins;
-    Hashset<std::pair<uint32_t, uint32_t>, 8> locationsAndIndexes;
+    Hashset<std::pair<uint32_t, uint32_t>, 8> locations_and_indices;
+    const ast::LocationAttribute* first_nonzero_location = nullptr;
+    const ast::IndexAttribute* first_nonzero_index = nullptr;
     enum class ParamOrRetType {
         kParameter,
         kReturnType,
@@ -1125,73 +1136,67 @@
         const ast::InterpolateAttribute* interpolate_attribute = nullptr;
         const ast::InvariantAttribute* invariant_attribute = nullptr;
         for (auto* attr : attrs) {
-            auto is_invalid_compute_shader_attribute = false;
+            bool ok = Switch(
+                attr,  //
+                [&](const ast::BuiltinAttribute* builtin_attr) {
+                    auto builtin = sem_.Get(builtin_attr)->Value();
 
-            if (auto* builtin_attr = attr->As<ast::BuiltinAttribute>()) {
-                auto builtin = sem_.Get(builtin_attr)->Value();
+                    if (pipeline_io_attribute) {
+                        AddError("multiple entry point IO attributes", attr->source);
+                        AddNote("previously consumed " + AttrToStr(pipeline_io_attribute),
+                                pipeline_io_attribute->source);
+                        return false;
+                    }
+                    pipeline_io_attribute = attr;
 
-                if (pipeline_io_attribute) {
-                    AddError("multiple entry point IO attributes", attr->source);
-                    AddNote("previously consumed " + AttrToStr(pipeline_io_attribute),
-                            pipeline_io_attribute->source);
-                    return false;
-                }
-                pipeline_io_attribute = attr;
+                    if (builtins.Contains(builtin)) {
+                        StringStream err;
+                        err << "@builtin(" << builtin << ") appears multiple times as pipeline "
+                            << (param_or_ret == ParamOrRetType::kParameter ? "input" : "output");
+                        AddError(err.str(), decl->source);
+                        return false;
+                    }
 
-                if (builtins.Contains(builtin)) {
-                    StringStream err;
-                    err << "@builtin(" << builtin << ") appears multiple times as pipeline "
-                        << (param_or_ret == ParamOrRetType::kParameter ? "input" : "output");
-                    AddError(err.str(), decl->source);
-                    return false;
-                }
+                    if (!BuiltinAttribute(
+                            builtin_attr, ty, stage,
+                            /* is_input */ param_or_ret == ParamOrRetType::kParameter)) {
+                        return false;
+                    }
+                    builtins.Add(builtin);
+                    return true;
+                },
+                [&](const ast::LocationAttribute* loc_attr) {
+                    location_attribute = loc_attr;
+                    if (pipeline_io_attribute) {
+                        AddError("multiple entry point IO attributes", attr->source);
+                        AddNote("previously consumed " + AttrToStr(pipeline_io_attribute),
+                                pipeline_io_attribute->source);
+                        return false;
+                    }
+                    pipeline_io_attribute = attr;
 
-                if (!BuiltinAttribute(builtin_attr, ty, stage,
-                                      /* is_input */ param_or_ret == ParamOrRetType::kParameter)) {
-                    return false;
-                }
-                builtins.Add(builtin);
-            } else if (auto* loc_attr = attr->As<ast::LocationAttribute>()) {
-                location_attribute = loc_attr;
-                if (pipeline_io_attribute) {
-                    AddError("multiple entry point IO attributes", attr->source);
-                    AddNote("previously consumed " + AttrToStr(pipeline_io_attribute),
-                            pipeline_io_attribute->source);
-                    return false;
-                }
-                pipeline_io_attribute = attr;
+                    if (TINT_UNLIKELY(!location.has_value())) {
+                        TINT_ICE() << "@location has no value";
+                        return false;
+                    }
 
-                bool is_input = param_or_ret == ParamOrRetType::kParameter;
+                    return LocationAttribute(loc_attr, ty, stage, source);
+                },
+                [&](const ast::IndexAttribute* index_attr) {
+                    index_attribute = index_attr;
+                    return IndexAttribute(index_attr, stage);
+                },
+                [&](const ast::InterpolateAttribute* interpolate) {
+                    interpolate_attribute = interpolate;
+                    return InterpolateAttribute(interpolate, ty, stage);
+                },
+                [&](const ast::InvariantAttribute* invariant) {
+                    invariant_attribute = invariant;
+                    return InvariantAttribute(invariant, stage);
+                },
+                [&](Default) { return true; });
 
-                if (TINT_UNLIKELY(!location.has_value())) {
-                    TINT_ICE() << "Location has no value";
-                    return false;
-                }
-
-                if (!LocationAttribute(loc_attr, ty, stage, source, is_input)) {
-                    return false;
-                }
-            } else if (auto* index_attr = attr->As<ast::IndexAttribute>()) {
-                index_attribute = index_attr;
-                return IndexAttribute(index_attr, stage);
-            } else if (auto* interpolate = attr->As<ast::InterpolateAttribute>()) {
-                if (decl->PipelineStage() == ast::PipelineStage::kCompute) {
-                    is_invalid_compute_shader_attribute = true;
-                } else if (!InterpolateAttribute(interpolate, ty)) {
-                    return false;
-                }
-                interpolate_attribute = interpolate;
-            } else if (auto* invariant = attr->As<ast::InvariantAttribute>()) {
-                if (decl->PipelineStage() == ast::PipelineStage::kCompute) {
-                    is_invalid_compute_shader_attribute = true;
-                }
-                invariant_attribute = invariant;
-            }
-            if (is_invalid_compute_shader_attribute) {
-                std::string input_or_output =
-                    param_or_ret == ParamOrRetType::kParameter ? "inputs" : "output";
-                AddError("@" + attr->Name() + " is not valid for compute shader " + input_or_output,
-                         attr->source);
+            if (!ok) {
                 return false;
             }
         }
@@ -1234,37 +1239,38 @@
             }
 
             if (index_attribute) {
-                if (Is<ast::LocationAttribute>(pipeline_io_attribute)) {
-                    AddError("index attribute must only be used with @location",
-                             index_attribute->source);
-                    return false;
-                }
-
                 // Because HLSL specifies dual source blending targets with SV_Target0 and 1, we
-                // should restrict targets with index attributes to location 0 for easy translation
+                // should restrict targets with @index to location 0 for easy translation
                 // in the backend writers.
-                if (location.value() != 0) {
-                    AddError("index attribute must only be used with @location(0)",
-                             index_attribute->source);
+                if (location.value_or(1) != 0) {
+                    AddError("@index can only be used with @location(0)", index_attribute->source);
                     return false;
                 }
             }
 
             if (location_attribute) {
-                uint32_t idx = 0xffffffff;
-                if (index_attribute) {
-                    idx = index.value();
+                if (!first_nonzero_location && location > 0u) {
+                    first_nonzero_location = location_attribute;
+                }
+                if (!first_nonzero_index && index > 0u) {
+                    first_nonzero_index = index_attribute;
+                }
+                if (first_nonzero_location && first_nonzero_index) {
+                    AddError("pipeline cannot use both non-zero @index and non-zero @location",
+                             first_nonzero_index->source);
+                    AddNote("non-zero @location declared here", first_nonzero_location->source);
+                    return false;
                 }
 
-                std::pair<uint32_t, uint32_t> locationAndIndex(location.value(), idx);
-                if (!locationsAndIndexes.Add(locationAndIndex)) {
+                std::pair<uint32_t, uint32_t> location_and_index(location.value(),
+                                                                 index.value_or(0));
+                if (!locations_and_indices.Add(location_and_index)) {
                     StringStream err;
-                    if (!index_attribute) {
-                        err << "@location(" << location.value() << ") appears multiple times";
-                    } else {
-                        err << "@location(" << location.value() << ") @index(" << index.value()
-                            << ") appears multiple times";
+                    err << "@location(" << location.value() << ") ";
+                    if (index_attribute) {
+                        err << "@index(" << index.value() << ") ";
                     }
+                    err << "appears multiple times";
                     AddError(err.str(), location_attribute->source);
                     return false;
                 }
@@ -1273,7 +1279,7 @@
             if (interpolate_attribute) {
                 if (!pipeline_io_attribute ||
                     !pipeline_io_attribute->Is<ast::LocationAttribute>()) {
-                    AddError("interpolate attribute must only be used with @location",
+                    AddError("@interpolate can only be used with @location",
                              interpolate_attribute->source);
                     return false;
                 }
@@ -1288,7 +1294,7 @@
                     }
                 }
                 if (!has_position) {
-                    AddError("invariant attribute must only be applied to a position builtin",
+                    AddError("@invariant must be applied to a position builtin",
                              invariant_attribute->source);
                     return false;
                 }
@@ -1327,9 +1333,10 @@
 
     for (auto* param : func->Parameters()) {
         auto* param_decl = param->Declaration();
+        auto& attrs = param->Attributes();
         if (!validate_entry_point_attributes(param_decl->attributes, param->Type(),
                                              param_decl->source, ParamOrRetType::kParameter,
-                                             param->Location(), std::nullopt)) {
+                                             attrs.location, attrs.index)) {
             return false;
         }
     }
@@ -1337,7 +1344,7 @@
     // Clear IO sets after parameter validation. Builtin and location attributes in return types
     // should be validated independently from those used in parameters.
     builtins.Clear();
-    locationsAndIndexes.Clear();
+    locations_and_indices.Clear();
 
     if (!func->ReturnType()->Is<core::type::Void>()) {
         if (!validate_entry_point_attributes(decl->return_type_attributes, func->ReturnType(),
@@ -1383,7 +1390,7 @@
         if (!var_decl) {
             continue;
         }
-        auto bp = global->BindingPoint();
+        auto bp = global->Attributes().binding_point;
         if (!bp) {
             continue;
         }
@@ -2146,8 +2153,7 @@
         return false;
     }
 
-    auto has_index = false;
-    Hashset<std::pair<uint32_t, uint32_t>, 8> locationsAndIndexes;
+    Hashset<std::pair<uint32_t, uint32_t>, 8> locations_and_indices;
     for (auto* member : str->Members()) {
         if (auto* r = member->Type()->As<sem::Array>()) {
             if (r->Count()->Is<core::type::RuntimeArrayCount>()) {
@@ -2179,20 +2185,16 @@
                 attr,  //
                 [&](const ast::InvariantAttribute* invariant) {
                     invariant_attribute = invariant;
-                    return true;
+                    return InvariantAttribute(invariant, stage);
                 },
                 [&](const ast::LocationAttribute* location) {
                     location_attribute = location;
                     TINT_ASSERT(member->Attributes().location.has_value());
-                    if (!LocationAttribute(location, member->Type(), stage,
-                                           member->Declaration()->source)) {
-                        return false;
-                    }
-                    return true;
+                    return LocationAttribute(location, member->Type(), stage,
+                                             member->Declaration()->source);
                 },
                 [&](const ast::IndexAttribute* index) {
                     index_attribute = index;
-                    has_index = true;
                     return IndexAttribute(index, stage);
                 },
                 [&](const ast::BuiltinAttribute* builtin_attr) {
@@ -2208,16 +2210,13 @@
                 },
                 [&](const ast::InterpolateAttribute* interpolate) {
                     interpolate_attribute = interpolate;
-                    if (!InterpolateAttribute(interpolate, member->Type())) {
-                        return false;
-                    }
-                    return true;
+                    return InterpolateAttribute(interpolate, member->Type(), stage);
                 },
                 [&](const ast::StructMemberSizeAttribute*) {
                     if (!member->Type()->HasCreationFixedFootprint()) {
                         AddError(
-                            "@size can only be applied to members where the member's type size "
-                            "can be fully determined at shader creation time",
+                            "@size can only be applied to members where the member's type size can "
+                            "be fully determined at shader creation time",
                             attr->source);
                         return false;
                     }
@@ -2230,62 +2229,40 @@
         }
 
         if (invariant_attribute && !has_position) {
-            AddError("invariant attribute must only be applied to a position builtin",
+            AddError("@invariant must be applied to a position builtin",
                      invariant_attribute->source);
             return false;
         }
 
         if (index_attribute) {
-            if (!location_attribute) {
-                AddError("index attribute must only be used with @location",
-                         index_attribute->source);
-                return false;
-            }
-
             // Because HLSL specifies dual source blending targets with SV_Target0 and 1, we should
             // restrict targets with index attributes to location 0 for easy translation in the
             // backend writers.
-            if (member->Attributes().location.value() != 0) {
-                AddError("index attribute must only be used with @location(0)",
-                         index_attribute->source);
+            if (member->Attributes().location.value_or(1) != 0) {
+                AddError("@index can only be used with @location(0)", index_attribute->source);
                 return false;
             }
         }
 
         if (interpolate_attribute && !location_attribute) {
-            AddError("interpolate attribute must only be used with @location",
-                     interpolate_attribute->source);
+            AddError("@interpolate can only be used with @location", interpolate_attribute->source);
             return false;
         }
 
         // Ensure all locations and index pairs are unique
         if (location_attribute) {
-            uint32_t index = 0xffffffff;
-            if (index_attribute) {
-                index = member->Attributes().index.value();
-            }
             uint32_t location = member->Attributes().location.value();
-            if (has_index && location != 0) {
-                StringStream err;
-                err << "Multiple render targets are not allowed when using dual source blending. "
-                       "The output @location("
-                    << location << ") is not allowed as a render target.";
-                AddError(err.str(), location_attribute->source);
-                return false;
-            }
+            uint32_t index = member->Attributes().index.value_or(0);
 
-            std::pair<uint32_t, uint32_t> locationAndIndex(location, index);
-            if (!locationsAndIndexes.Add(locationAndIndex)) {
+            std::pair<uint32_t, uint32_t> location_and_index(location, index);
+            if (!locations_and_indices.Add(location_and_index)) {
                 StringStream err;
-                if (!index_attribute) {
-                    err << "@location(" << location << ") appears multiple times";
-                    AddError(err.str(), location_attribute->source);
-                } else {
-                    err << "@location(" << location << ") @index(" << index
-                        << ") appears multiple times";
-                    AddError(err.str(), index_attribute->source);
+                err << "@location(" << location << ") ";
+                if (index_attribute) {
+                    err << "@index(" << index << ") ";
                 }
-
+                err << "appears multiple times";
+                AddError(err.str(), location_attribute->source);
                 return false;
             }
         }
@@ -2294,15 +2271,12 @@
     return true;
 }
 
-bool Validator::LocationAttribute(const ast::LocationAttribute* loc_attr,
+bool Validator::LocationAttribute(const ast::LocationAttribute* attr,
                                   const core::type::Type* type,
                                   ast::PipelineStage stage,
-                                  const Source& source,
-                                  const bool is_input) const {
-    std::string inputs_or_output = is_input ? "inputs" : "output";
+                                  const Source& source) const {
     if (stage == ast::PipelineStage::kCompute) {
-        AddError("@" + loc_attr->Name() + " is not valid for compute shader " + inputs_or_output,
-                 loc_attr->source);
+        AddError(AttrToStr(attr) + " cannot be used by compute shaders", attr->source);
         return false;
     }
 
@@ -2312,32 +2286,28 @@
         AddNote(
             "@location must only be applied to declarations of numeric scalar or numeric vector "
             "type",
-            loc_attr->source);
+            attr->source);
         return false;
     }
 
     return true;
 }
 
-bool Validator::IndexAttribute(const ast::IndexAttribute* index_attr,
-                               ast::PipelineStage stage) const {
+bool Validator::IndexAttribute(const ast::IndexAttribute* attr,
+                               ast::PipelineStage stage,
+                               const std::optional<bool> is_input) const {
     if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumInternalDualSourceBlending)) {
         AddError(
-            "use of '@index' attribute requires enabling extension "
-            "'chromium_internal_dual_source_blending'",
-            index_attr->source);
+            "use of @index requires enabling extension 'chromium_internal_dual_source_blending'",
+            attr->source);
         return false;
     }
 
-    if (stage == ast::PipelineStage::kCompute) {
-        AddError("@" + index_attr->Name() + " is not valid for compute shader output",
-                 index_attr->source);
-        return false;
-    }
-
-    if (stage == ast::PipelineStage::kVertex) {
-        AddError("@" + index_attr->Name() + " is not valid for vertex shader output",
-                 index_attr->source);
+    bool is_stage_non_fragment =
+        stage != ast::PipelineStage::kNone && stage != ast::PipelineStage::kFragment;
+    bool is_output = is_input.value_or(false);
+    if (is_stage_non_fragment || is_output) {
+        AddError(AttrToStr(attr) + " can only be used for fragment shader output", attr->source);
         return false;
     }
 
diff --git a/src/tint/lang/wgsl/resolver/validator.h b/src/tint/lang/wgsl/resolver/validator.h
index 20735de..03466ee 100644
--- a/src/tint/lang/wgsl/resolver/validator.h
+++ b/src/tint/lang/wgsl/resolver/validator.h
@@ -315,11 +315,20 @@
     bool IncrementDecrementStatement(const ast::IncrementDecrementStatement* stmt) const;
 
     /// Validates an interpolate attribute
-    /// @param attr the interpolation attribute to validate
+    /// @param attr the attribute to validate
     /// @param storage_type the storage type of the attached variable
-    /// @returns true on succes, false otherwise
+    /// @param stage the current pipeline stage
+    /// @returns true on success, false otherwise
     bool InterpolateAttribute(const ast::InterpolateAttribute* attr,
-                              const core::type::Type* storage_type) const;
+                              const core::type::Type* storage_type,
+                              const ast::PipelineStage stage) const;
+
+    /// Validates an invariant attribute
+    /// @param attr the attribute to validate
+    /// @param stage the current pipeline stage
+    /// @returns true on success, false otherwise
+    bool InvariantAttribute(const ast::InvariantAttribute* attr,
+                            const ast::PipelineStage stage) const;
 
     /// Validates a builtin call
     /// @param call the builtin call to validate
@@ -332,23 +341,25 @@
     bool LocalVariable(const sem::Variable* v) const;
 
     /// Validates a location attribute
-    /// @param loc_attr the location attribute to validate
+    /// @param attr the attribute to validate
     /// @param type the variable type
     /// @param stage the current pipeline stage
-    /// @param source the source of the attribute
-    /// @param is_input true if this is an input variable
+    /// @param source the source of declaration using the attribute
     /// @returns true on success, false otherwise.
-    bool LocationAttribute(const ast::LocationAttribute* loc_attr,
+    bool LocationAttribute(const ast::LocationAttribute* attr,
                            const core::type::Type* type,
-                           ast::PipelineStage stage,
-                           const Source& source,
-                           const bool is_input = false) const;
+                           const ast::PipelineStage stage,
+                           const Source& source) const;
 
     /// Validates a index attribute
     /// @param index_attr the index attribute to validate
     /// @param stage the current pipeline stage
+    /// @param is_input true if is an input variable, false if output variable, std::nullopt is
+    /// unknown.
     /// @returns true on success, false otherwise.
-    bool IndexAttribute(const ast::IndexAttribute* index_attr, ast::PipelineStage stage) const;
+    bool IndexAttribute(const ast::IndexAttribute* index_attr,
+                        ast::PipelineStage stage,
+                        const std::optional<bool> is_input = std::nullopt) const;
 
     /// Validates a loop statement
     /// @param stmt the loop statement
diff --git a/src/tint/lang/wgsl/sem/function.cc b/src/tint/lang/wgsl/sem/function.cc
index 4d1d3c9..9a56376 100644
--- a/src/tint/lang/wgsl/sem/function.cc
+++ b/src/tint/lang/wgsl/sem/function.cc
@@ -81,7 +81,7 @@
             continue;
         }
 
-        if (auto bp = global->BindingPoint()) {
+        if (auto bp = global->Attributes().binding_point) {
             ret.push_back({global, *bp});
         }
     }
@@ -96,7 +96,7 @@
             continue;
         }
 
-        if (auto bp = global->BindingPoint()) {
+        if (auto bp = global->Attributes().binding_point) {
             ret.push_back({global, *bp});
         }
     }
@@ -140,7 +140,7 @@
     for (auto* global : TransitivelyReferencedGlobals()) {
         auto* unwrapped_type = global->Type()->UnwrapRef();
         if (unwrapped_type->TypeInfo().Is(type)) {
-            if (auto bp = global->BindingPoint()) {
+            if (auto bp = global->Attributes().binding_point) {
                 ret.push_back({global, *bp});
             }
         }
@@ -168,7 +168,7 @@
             continue;
         }
 
-        if (auto bp = global->BindingPoint()) {
+        if (auto bp = global->Attributes().binding_point) {
             ret.push_back({global, *bp});
         }
     }
@@ -193,7 +193,7 @@
             continue;
         }
 
-        if (auto bp = global->BindingPoint()) {
+        if (auto bp = global->Attributes().binding_point) {
             ret.push_back({global, *bp});
         }
     }
diff --git a/src/tint/lang/wgsl/sem/variable.h b/src/tint/lang/wgsl/sem/variable.h
index bca5d1b..3ee24f3 100644
--- a/src/tint/lang/wgsl/sem/variable.h
+++ b/src/tint/lang/wgsl/sem/variable.h
@@ -151,6 +151,22 @@
     const CastableBase* shadows_ = nullptr;
 };
 
+/// Attributes that can be applied to global variables
+struct GlobalVariableAttributes {
+    /// the pipeline constant ID associated with the variable
+    std::optional<tint::OverrideId> override_id;
+    /// the resource binding point for the variable, if set.
+    std::optional<tint::BindingPoint> binding_point;
+    /// The `location` attribute value for the variable, if set
+    /// @note a GlobalVariable generally doesn't have a `location` in WGSL, as it isn't allowed by
+    /// the spec. The location maybe attached by transforms such as CanonicalizeEntryPointIO.
+    std::optional<uint32_t> location;
+    /// The `index` attribute value for the variable, if set
+    /// @note a GlobalVariable generally doesn't have a `index` in WGSL, as it isn't allowed by
+    /// the spec. The location maybe attached by transforms such as CanonicalizeEntryPointIO.
+    std::optional<uint32_t> index;
+};
+
 /// GlobalVariable is a module-scope variable
 class GlobalVariable final : public Castable<GlobalVariable, Variable> {
   public:
@@ -161,34 +177,6 @@
     /// Destructor
     ~GlobalVariable() override;
 
-    /// @param binding_point the resource binding point for the parameter
-    void SetBindingPoint(std::optional<tint::BindingPoint> binding_point) {
-        binding_point_ = binding_point;
-    }
-
-    /// @returns the resource binding point for the variable
-    std::optional<tint::BindingPoint> BindingPoint() const { return binding_point_; }
-
-    /// @param id the constant identifier to assign to this variable
-    void SetOverrideId(OverrideId id) { override_id_ = id; }
-
-    /// @returns the pipeline constant ID associated with the variable
-    tint::OverrideId OverrideId() const { return override_id_; }
-
-    /// @param location the location value for the parameter, if set
-    /// @note a GlobalVariable generally doesn't have a `location` in WGSL, as it isn't allowed by
-    /// the spec. The location maybe attached by transforms such as CanonicalizeEntryPointIO.
-    void SetLocation(std::optional<uint32_t> location) { location_ = location; }
-
-    /// @returns the location value for the parameter, if set
-    std::optional<uint32_t> Location() const { return location_; }
-
-    /// @param index the index value for the parameter, if set
-    void SetIndex(std::optional<uint32_t> index) { index_ = index; }
-
-    /// @returns the index value for the parameter, if set
-    std::optional<uint32_t> Index() const { return index_; }
-
     /// Records that this variable (transitively) references the given override variable.
     /// @param var the module-scope override variable
     void AddTransitivelyReferencedOverride(const GlobalVariable* var);
@@ -198,12 +186,30 @@
         return transitively_referenced_overrides_;
     }
 
+    /// @return the mutable attributes for the variable
+    GlobalVariableAttributes& Attributes() { return attributes_; }
+
+    /// @return the immutable attributes for the variable
+    const GlobalVariableAttributes& Attributes() const { return attributes_; }
+
   private:
     std::optional<tint::BindingPoint> binding_point_;
     tint::OverrideId override_id_;
-    std::optional<uint32_t> location_;
-    std::optional<uint32_t> index_;
     UniqueVector<const GlobalVariable*, 4> transitively_referenced_overrides_;
+    GlobalVariableAttributes attributes_;
+};
+
+/// Attributes that can be applied to parameters
+struct ParameterAttributes {
+    /// the resource binding point for the variable, if set.
+    /// @note a Parameter generally doesn't have a `group` or `binding` attribute in WGSL, as it
+    /// isn't allowed by the spec. The binding point maybe attached by transforms such as
+    /// CanonicalizeEntryPointIO.
+    std::optional<tint::BindingPoint> binding_point;
+    /// The `location` attribute value for the variable, if set
+    std::optional<uint32_t> location;
+    /// The `index` attribute value for the variable, if set
+    std::optional<uint32_t> index;
 };
 
 /// Parameter is a function parameter
@@ -227,9 +233,6 @@
         return static_cast<const ast::Parameter*>(Variable::Declaration());
     }
 
-    /// @param index the index value for the parameter, if set
-    void SetIndex(uint32_t index) { index_ = index; }
-
     /// @return the index of the parameter in the function
     uint32_t Index() const { return index_; }
 
@@ -252,27 +255,18 @@
     /// @returns the Type, Function or Variable that this local variable shadows
     const CastableBase* Shadows() const { return shadows_; }
 
-    /// @param binding_point the resource binding point for the parameter
-    void SetBindingPoint(std::optional<tint::BindingPoint> binding_point) {
-        binding_point_ = binding_point;
-    }
+    /// @return the mutable attributes for the parameter
+    ParameterAttributes& Attributes() { return attributes_; }
 
-    /// @returns the resource binding point for the parameter
-    std::optional<tint::BindingPoint> BindingPoint() const { return binding_point_; }
-
-    /// @param location the location value for the parameter, if set
-    void SetLocation(std::optional<uint32_t> location) { location_ = location; }
-
-    /// @returns the location value for the parameter, if set
-    std::optional<uint32_t> Location() const { return location_; }
+    /// @return the immutable attributes for the parameter
+    const ParameterAttributes& Attributes() const { return attributes_; }
 
   private:
-    uint32_t index_ = 0;
+    const uint32_t index_ = 0;
     core::ParameterUsage usage_ = core::ParameterUsage::kNone;
     CallTarget const* owner_ = nullptr;
     const CastableBase* shadows_ = nullptr;
-    std::optional<tint::BindingPoint> binding_point_;
-    std::optional<uint32_t> location_;
+    ParameterAttributes attributes_;
 };
 
 /// VariableUser holds the semantic information for an identifier expression
diff --git a/src/tint/lang/wgsl/wgsl.def b/src/tint/lang/wgsl/wgsl.def
index a66c2ba..8c545c9 100644
--- a/src/tint/lang/wgsl/wgsl.def
+++ b/src/tint/lang/wgsl/wgsl.def
@@ -87,7 +87,11 @@
   // A Chromium-specific extension that enables dual source blending.
   chromium_internal_dual_source_blending
   // A Chromium-specific extension that enables pixel local storage.
+  // Cannot be used with chromium_experimental_framebuffer_fetch
   chromium_experimental_pixel_local
+  // A Chromium-specific extension that enables framebuffer fetching.
+  // Cannot be used with chromium_experimental_pixel_local
+  chromium_experimental_framebuffer_fetch
 }
 
 // https://gpuweb.github.io/gpuweb/wgsl/#language-extensions-sec