tint/resolver: Resolve 'const' variables

The frontends do not currently emit these, nor do the backends currently
handle them.

Bug: tint:1580
Change-Id: I469a5379663c2802145b28a94f5c1e348cc14ff3
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/94605
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>

src/tint/program_builder.h

@@ -33,6 +33,7 @@
 #include "src/tint/ast/call_statement.h"
 #include "src/tint/ast/case_statement.h"
 #include "src/tint/ast/compound_assignment_statement.h"
+#include "src/tint/ast/const.h"
 #include "src/tint/ast/continue_statement.h"
 #include "src/tint/ast/depth_multisampled_texture.h"
 #include "src/tint/ast/depth_texture.h"
@@ -1363,6 +1364,35 @@ class ProgramBuilder {
                                 opts.access, opts.constructor, std::move(opts.attributes));
     }
 
+    /// @param name the variable name
+    /// @param type the variable type
+    /// @param constructor constructor expression
+    /// @param attributes optional variable attributes
+    /// @returns an `ast::Const` with the given name and type
+    template <typename NAME>
+    const ast::Const* Const(NAME&& name,
+                            const ast::Type* type,
+                            const ast::Expression* constructor,
+                            ast::AttributeList attributes = {}) {
+        return create<ast::Const>(Sym(std::forward<NAME>(name)), type, constructor, attributes);
+    }
+
+    /// @param source the variable source
+    /// @param name the variable name
+    /// @param type the variable type
+    /// @param constructor constructor expression
+    /// @param attributes optional variable attributes
+    /// @returns an `ast::Const` with the given name and type
+    template <typename NAME>
+    const ast::Const* Const(const Source& source,
+                            NAME&& name,
+                            const ast::Type* type,
+                            const ast::Expression* constructor,
+                            ast::AttributeList attributes = {}) {
+        return create<ast::Const>(source, Sym(std::forward<NAME>(name)), type, constructor,
+                                  attributes);
+    }
+
     /// @param name the variable name
     /// @param type the variable type
     /// @param constructor constructor expression
@@ -1459,6 +1489,42 @@ class ProgramBuilder {
         return var;
     }
 
+    /// @param name the variable name
+    /// @param type the variable type
+    /// @param constructor constructor expression
+    /// @param attributes optional variable attributes
+    /// @returns an `ast::Const` constructed by calling Const() with the arguments of `args`, which
+    /// is automatically registered as a global variable with the ast::Module.
+    template <typename NAME>
+    const ast::Const* GlobalConst(NAME&& name,
+                                  const ast::Type* type,
+                                  const ast::Expression* constructor,
+                                  ast::AttributeList attributes = {}) {
+        auto* var = Const(std::forward<NAME>(name), type, constructor, std::move(attributes));
+        AST().AddGlobalVariable(var);
+        return var;
+    }
+
+    /// @param source the variable source
+    /// @param name the variable name
+    /// @param type the variable type
+    /// @param constructor constructor expression
+    /// @param attributes optional variable attributes
+    /// @returns a const `ast::Const` constructed by calling Var() with the
+    /// arguments of `args`, which is automatically registered as a global
+    /// variable with the ast::Module.
+    template <typename NAME>
+    const ast::Const* GlobalConst(const Source& source,
+                                  NAME&& name,
+                                  const ast::Type* type,
+                                  const ast::Expression* constructor,
+                                  ast::AttributeList attributes = {}) {
+        auto* var =
+            Const(source, std::forward<NAME>(name), type, constructor, std::move(attributes));
+        AST().AddGlobalVariable(var);
+        return var;
+    }
+
     /// @param name the variable name
     /// @param type the variable type
     /// @param constructor constructor expression

src/tint/resolver/array_accessor_test.cc

@@ -65,7 +65,7 @@ TEST_F(ResolverIndexAccessorTest, Matrix_BothDimensions_Dynamic_Ref) {
 }
 
 TEST_F(ResolverIndexAccessorTest, Matrix_Dynamic) {
-    GlobalLet("my_const", ty.mat2x3<f32>(), Construct(ty.mat2x3<f32>()));
+    GlobalConst("my_const", ty.mat2x3<f32>(), Construct(ty.mat2x3<f32>()));
     auto* idx = Var("idx", ty.i32(), Construct(ty.i32()));
     auto* acc = IndexAccessor("my_const", Expr(Source{{12, 34}}, idx));
     WrapInFunction(Decl(idx), acc);
@@ -79,10 +79,32 @@ TEST_F(ResolverIndexAccessorTest, Matrix_Dynamic) {
     EXPECT_EQ(idx_sem->Object()->Declaration(), acc->object);
 }
 
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverIndexAccessorTest, Matrix_Dynamic_Let) {
+    GlobalLet("my_let", ty.mat2x3<f32>(), Construct(ty.mat2x3<f32>()));
+    auto* idx = Var("idx", ty.i32(), Construct(ty.i32()));
+    auto* acc = IndexAccessor("my_let", Expr(Source{{12, 34}}, idx));
+    WrapInFunction(Decl(idx), acc);
+
+    EXPECT_TRUE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "");
+}
+
 TEST_F(ResolverIndexAccessorTest, Matrix_XDimension_Dynamic) {
-    GlobalLet("my_var", ty.mat4x4<f32>(), Construct(ty.mat4x4<f32>()));
+    GlobalConst("my_const", ty.mat4x4<f32>(), Construct(ty.mat4x4<f32>()));
     auto* idx = Var("idx", ty.u32(), Expr(3_u));
-    auto* acc = IndexAccessor("my_var", Expr(Source{{12, 34}}, idx));
+    auto* acc = IndexAccessor("my_const", Expr(Source{{12, 34}}, idx));
+    WrapInFunction(Decl(idx), acc);
+
+    EXPECT_TRUE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverIndexAccessorTest, Matrix_XDimension_Dynamic_Let) {
+    GlobalLet("my_let", ty.mat4x4<f32>(), Construct(ty.mat4x4<f32>()));
+    auto* idx = Var("idx", ty.u32(), Expr(3_u));
+    auto* acc = IndexAccessor("my_let", Expr(Source{{12, 34}}, idx));
     WrapInFunction(Decl(idx), acc);
 
     EXPECT_TRUE(r()->Resolve());
@@ -95,9 +117,20 @@ TEST_F(ResolverIndexAccessorTest, Matrix_XDimension_Dynamic) {
 }
 
 TEST_F(ResolverIndexAccessorTest, Matrix_BothDimension_Dynamic) {
-    GlobalLet("my_var", ty.mat4x4<f32>(), Construct(ty.mat4x4<f32>()));
+    GlobalConst("my_const", ty.mat4x4<f32>(), Construct(ty.mat4x4<f32>()));
     auto* idx = Var("idy", ty.u32(), Expr(2_u));
-    auto* acc = IndexAccessor(IndexAccessor("my_var", Expr(Source{{12, 34}}, idx)), 1_i);
+    auto* acc = IndexAccessor(IndexAccessor("my_const", Expr(Source{{12, 34}}, idx)), 1_i);
+    WrapInFunction(Decl(idx), acc);
+
+    EXPECT_TRUE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverIndexAccessorTest, Matrix_BothDimension_Dynamic_Let) {
+    GlobalLet("my_let", ty.mat4x4<f32>(), Construct(ty.mat4x4<f32>()));
+    auto* idx = Var("idy", ty.u32(), Expr(2_u));
+    auto* acc = IndexAccessor(IndexAccessor("my_let", Expr(Source{{12, 34}}, idx)), 1_i);
     WrapInFunction(Decl(idx), acc);
 
     EXPECT_TRUE(r()->Resolve());
@@ -174,9 +207,19 @@ TEST_F(ResolverIndexAccessorTest, Vector_Dynamic_Ref) {
 }
 
 TEST_F(ResolverIndexAccessorTest, Vector_Dynamic) {
-    GlobalLet("my_var", ty.vec3<f32>(), Construct(ty.vec3<f32>()));
+    GlobalConst("my_const", ty.vec3<f32>(), Construct(ty.vec3<f32>()));
     auto* idx = Var("idx", ty.i32(), Expr(2_i));
-    auto* acc = IndexAccessor("my_var", Expr(Source{{12, 34}}, idx));
+    auto* acc = IndexAccessor("my_const", Expr(Source{{12, 34}}, idx));
+    WrapInFunction(Decl(idx), acc);
+
+    EXPECT_TRUE(r()->Resolve());
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverIndexAccessorTest, Vector_Dynamic_Let) {
+    GlobalLet("my_let", ty.vec3<f32>(), Construct(ty.vec3<f32>()));
+    auto* idx = Var("idx", ty.i32(), Expr(2_i));
+    auto* acc = IndexAccessor("my_let", Expr(Source{{12, 34}}, idx));
     WrapInFunction(Decl(idx), acc);
 
     EXPECT_TRUE(r()->Resolve());
@@ -278,9 +321,22 @@ TEST_F(ResolverIndexAccessorTest, Alias_Array) {
 }
 
 TEST_F(ResolverIndexAccessorTest, Array_Constant) {
-    GlobalLet("my_var", ty.array<f32, 3>(), array<f32, 3>());
+    GlobalConst("my_const", ty.array<f32, 3>(), array<f32, 3>());
 
-    auto* acc = IndexAccessor("my_var", 2_i);
+    auto* acc = IndexAccessor("my_const", 2_i);
+    WrapInFunction(acc);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    ASSERT_NE(TypeOf(acc), nullptr);
+    EXPECT_TRUE(TypeOf(acc)->Is<sem::F32>());
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverIndexAccessorTest, Array_Let) {
+    GlobalLet("my_let", ty.array<f32, 3>(), array<f32, 3>());
+
+    auto* acc = IndexAccessor("my_let", 2_i);
     WrapInFunction(acc);
 
     EXPECT_TRUE(r()->Resolve()) << r()->error();

src/tint/resolver/assignment_validation_test.cc

@@ -61,6 +61,26 @@ TEST_F(ResolverAssignmentValidationTest, AssignIncompatibleTypes) {
 }
 
 TEST_F(ResolverAssignmentValidationTest, AssignArraysWithDifferentSizeExpressions_Pass) {
+    // const len = 4u;
+    // {
+    //   var a : array<f32, 4u>;
+    //   var b : array<f32, len>;
+    //   a = b;
+    // }
+
+    GlobalConst("len", nullptr, Expr(4_u));
+
+    auto* a = Var("a", ty.array(ty.f32(), 4_u));
+    auto* b = Var("b", ty.array(ty.f32(), "len"));
+
+    auto* assign = Assign(Source{{12, 34}}, "a", "b");
+    WrapInFunction(a, b, assign);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverAssignmentValidationTest, AssignArraysWithDifferentSizeExpressions_Let_Pass) {
     // let len = 4u;
     // {
     //   var a : array<f32, 4u>;
@@ -80,6 +100,28 @@ TEST_F(ResolverAssignmentValidationTest, AssignArraysWithDifferentSizeExpression
 }
 
 TEST_F(ResolverAssignmentValidationTest, AssignArraysWithDifferentSizeExpressions_Fail) {
+    // const len = 5u;
+    // {
+    //   var a : array<f32, 4u>;
+    //   var b : array<f32, len>;
+    //   a = b;
+    // }
+
+    GlobalConst("len", nullptr, Expr(5_u));
+
+    auto* a = Var("a", ty.array(ty.f32(), 4_u));
+    auto* b = Var("b", ty.array(ty.f32(), "len"));
+
+    auto* assign = Assign(Source{{12, 34}}, "a", "b");
+    WrapInFunction(a, b, assign);
+
+    ASSERT_FALSE(r()->Resolve());
+
+    EXPECT_EQ(r()->error(), "12:34 error: cannot assign 'array<f32, 5>' to 'array<f32, 4>'");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverAssignmentValidationTest, AssignArraysWithDifferentSizeExpressions_Let_Fail) {
     // let len = 5u;
     // {
     //   var a : array<f32, 4u>;

src/tint/resolver/attribute_validation_test.cc

@@ -774,6 +774,24 @@ using ConstantAttributeTest = TestWithParams;
 TEST_P(ConstantAttributeTest, IsValid) {
     auto& params = GetParam();
 
+    GlobalConst("a", ty.f32(), Expr(1.23_f),
+                createAttributes(Source{{12, 34}}, *this, params.kind));
+
+    WrapInFunction();
+
+    if (params.should_pass) {
+        EXPECT_TRUE(r()->Resolve()) << r()->error();
+    } else {
+        EXPECT_FALSE(r()->Resolve());
+        EXPECT_EQ(r()->error(),
+                  "12:34 error: attribute is not valid for module-scope 'const' declaration");
+    }
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_P(ConstantAttributeTest, IsValid_Let) {
+    auto& params = GetParam();
+
     GlobalLet("a", ty.f32(), Expr(1.23_f), createAttributes(Source{{12, 34}}, *this, params.kind));
 
     WrapInFunction();
@@ -804,6 +822,22 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kBindingAndGroup, false}));
 
 TEST_F(ConstantAttributeTest, DuplicateAttribute) {
+    GlobalConst("a", ty.f32(), Expr(1.23_f),
+                ast::AttributeList{
+                    create<ast::IdAttribute>(Source{{12, 34}}, 0),
+                    create<ast::IdAttribute>(Source{{56, 78}}, 1),
+                });
+
+    WrapInFunction();
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              R"(56:78 error: duplicate id attribute
+12:34 note: first attribute declared here)");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ConstantAttributeTest, DuplicateAttribute_Let) {
     GlobalLet("a", ty.f32(), Expr(1.23_f),
               ast::AttributeList{
                   create<ast::IdAttribute>(Source{{12, 34}}, 0),
@@ -851,7 +885,7 @@ INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
                                          TestParams{AttributeKind::kBindingAndGroup, false}));
 
 TEST_F(OverrideAttributeTest, DuplicateAttribute) {
-    GlobalLet("a", ty.f32(), Expr(1.23_f),
+    Override("a", ty.f32(), Expr(1.23_f),
              ast::AttributeList{
                  create<ast::IdAttribute>(Source{{12, 34}}, 0),
                  create<ast::IdAttribute>(Source{{56, 78}}, 1),

src/tint/resolver/builtin_validation_test.cc

@@ -85,6 +85,15 @@ TEST_F(ResolverBuiltinValidationTest, BuiltinRedeclaredAsFunction) {
               R"(12:34 error: 'mix' is a builtin and cannot be redeclared as a function)");
 }
 
+TEST_F(ResolverBuiltinValidationTest, BuiltinRedeclaredAsGlobalConst) {
+    GlobalConst(Source{{12, 34}}, "mix", ty.i32(), Expr(1_i));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              R"(12:34 error: 'mix' is a builtin and cannot be redeclared as a 'const')");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverBuiltinValidationTest, BuiltinRedeclaredAsGlobalLet) {
     GlobalLet(Source{{12, 34}}, "mix", ty.i32(), Expr(1_i));
 
@@ -267,6 +276,39 @@ TEST_P(BuiltinTextureConstExprArgValidationTest, Immediate) {
     }
 }
 
+TEST_P(BuiltinTextureConstExprArgValidationTest, GlobalConst) {
+    auto& p = GetParam();
+    auto overload = std::get<0>(p);
+    auto param = std::get<1>(p);
+    auto expr = std::get<2>(p);
+
+    // Build the global texture and sampler variables
+    overload.BuildTextureVariable(this);
+    overload.BuildSamplerVariable(this);
+
+    // Build the module-scope const 'G' with the offset value
+    GlobalConst("G", nullptr, expr({}, *this));
+
+    auto args = overload.args(this);
+    auto*& arg_to_replace = (param.position == Position::kFirst) ? args.front() : args.back();
+
+    // Make the expression to be replaced, reachable. This keeps the resolver
+    // happy.
+    WrapInFunction(arg_to_replace);
+
+    arg_to_replace = Expr(Source{{12, 34}}, "G");
+
+    // Call the builtin with the constexpr argument replaced
+    Func("func", {}, ty.void_(), {CallStmt(Call(overload.function, args))},
+         {Stage(ast::PipelineStage::kFragment)});
+
+    EXPECT_FALSE(r()->Resolve());
+    std::stringstream err;
+    err << "12:34 error: the " << param.name << " argument must be a const_expression";
+    EXPECT_EQ(r()->error(), err.str());
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_P(BuiltinTextureConstExprArgValidationTest, GlobalLet) {
     auto& p = GetParam();
     auto overload = std::get<0>(p);

src/tint/resolver/dependency_graph.cc

@@ -491,9 +491,10 @@ struct DependencyAnalysis {
             [&](const ast::Struct*) { return "struct"; },      //
             [&](const ast::Alias*) { return "alias"; },        //
             [&](const ast::Function*) { return "function"; },  //
-            [&](const ast::Let*) { return "let"; },            //
             [&](const ast::Var*) { return "var"; },            //
+            [&](const ast::Let*) { return "let"; },            //
             [&](const ast::Override*) { return "override"; },  //
+            [&](const ast::Const*) { return "const"; },        //
             [&](Default) {
                 UnhandledNode(diagnostics_, node);
                 return "<error>";

src/tint/resolver/dependency_graph_test.cc

@@ -55,6 +55,7 @@ using ResolverDependencyGraphTest = ResolverDependencyGraphTestWithParam<::testi
 enum class SymbolDeclKind {
     GlobalVar,
     GlobalLet,
+    GlobalConst,
     Alias,
     Struct,
     Function,
@@ -66,10 +67,10 @@ enum class SymbolDeclKind {
 };
 
 static constexpr SymbolDeclKind kAllSymbolDeclKinds[] = {
-    SymbolDeclKind::GlobalVar,      SymbolDeclKind::GlobalLet, SymbolDeclKind::Alias,
+    SymbolDeclKind::GlobalVar,      SymbolDeclKind::GlobalLet,      SymbolDeclKind::GlobalConst,
-    SymbolDeclKind::Struct,         SymbolDeclKind::Function,  SymbolDeclKind::Parameter,
+    SymbolDeclKind::Alias,          SymbolDeclKind::Struct,         SymbolDeclKind::Function,
-    SymbolDeclKind::LocalVar,       SymbolDeclKind::LocalLet,  SymbolDeclKind::NestedLocalVar,
+    SymbolDeclKind::Parameter,      SymbolDeclKind::LocalVar,       SymbolDeclKind::LocalLet,
-    SymbolDeclKind::NestedLocalLet,
+    SymbolDeclKind::NestedLocalVar, SymbolDeclKind::NestedLocalLet,
 };
 
 static constexpr SymbolDeclKind kTypeDeclKinds[] = {
@@ -78,14 +79,14 @@ static constexpr SymbolDeclKind kTypeDeclKinds[] = {
 };
 
 static constexpr SymbolDeclKind kValueDeclKinds[] = {
-    SymbolDeclKind::GlobalVar,      SymbolDeclKind::GlobalLet, SymbolDeclKind::Parameter,
+    SymbolDeclKind::GlobalVar,      SymbolDeclKind::GlobalLet,      SymbolDeclKind::GlobalConst,
-    SymbolDeclKind::LocalVar,       SymbolDeclKind::LocalLet,  SymbolDeclKind::NestedLocalVar,
+    SymbolDeclKind::Parameter,      SymbolDeclKind::LocalVar,       SymbolDeclKind::LocalLet,
-    SymbolDeclKind::NestedLocalLet,
+    SymbolDeclKind::NestedLocalVar, SymbolDeclKind::NestedLocalLet,
 };
 
 static constexpr SymbolDeclKind kGlobalDeclKinds[] = {
-    SymbolDeclKind::GlobalVar, SymbolDeclKind::GlobalLet, SymbolDeclKind::Alias,
+    SymbolDeclKind::GlobalVar, SymbolDeclKind::GlobalLet, SymbolDeclKind::GlobalConst,
-    SymbolDeclKind::Struct,    SymbolDeclKind::Function,
+    SymbolDeclKind::Alias,     SymbolDeclKind::Struct,    SymbolDeclKind::Function,
 };
 
 static constexpr SymbolDeclKind kLocalDeclKinds[] = {
@@ -96,6 +97,7 @@ static constexpr SymbolDeclKind kLocalDeclKinds[] = {
 static constexpr SymbolDeclKind kGlobalValueDeclKinds[] = {
     SymbolDeclKind::GlobalVar,
     SymbolDeclKind::GlobalLet,
+    SymbolDeclKind::GlobalConst,
 };
 
 static constexpr SymbolDeclKind kFuncDeclKinds[] = {
@@ -119,6 +121,12 @@ enum class SymbolUseKind {
     GlobalLetVectorElemType,
     GlobalLetMatrixElemType,
     GlobalLetValue,
+    GlobalConstType,
+    GlobalConstArrayElemType,
+    GlobalConstArraySizeValue,
+    GlobalConstVectorElemType,
+    GlobalConstMatrixElemType,
+    GlobalConstValue,
     AliasType,
     StructMemberType,
     CallFunction,
@@ -151,6 +159,11 @@ static constexpr SymbolUseKind kTypeUseKinds[] = {
     SymbolUseKind::GlobalLetArraySizeValue,
     SymbolUseKind::GlobalLetVectorElemType,
     SymbolUseKind::GlobalLetMatrixElemType,
+    SymbolUseKind::GlobalConstType,
+    SymbolUseKind::GlobalConstArrayElemType,
+    SymbolUseKind::GlobalConstArraySizeValue,
+    SymbolUseKind::GlobalConstVectorElemType,
+    SymbolUseKind::GlobalConstMatrixElemType,
     SymbolUseKind::AliasType,
     SymbolUseKind::StructMemberType,
     SymbolUseKind::ParameterType,
@@ -166,9 +179,9 @@ static constexpr SymbolUseKind kTypeUseKinds[] = {
 
 static constexpr SymbolUseKind kValueUseKinds[] = {
     SymbolUseKind::GlobalVarValue,      SymbolUseKind::GlobalLetValue,
-    SymbolUseKind::LocalVarValue,       SymbolUseKind::LocalLetValue,
+    SymbolUseKind::GlobalConstValue,    SymbolUseKind::LocalVarValue,
-    SymbolUseKind::NestedLocalVarValue, SymbolUseKind::NestedLocalLetValue,
+    SymbolUseKind::LocalLetValue,       SymbolUseKind::NestedLocalVarValue,
-    SymbolUseKind::WorkgroupSizeValue,
+    SymbolUseKind::NestedLocalLetValue, SymbolUseKind::WorkgroupSizeValue,
 };
 
 static constexpr SymbolUseKind kFuncUseKinds[] = {
@@ -183,6 +196,8 @@ std::ostream& operator<<(std::ostream& out, SymbolDeclKind kind) {
             return out << "global var";
         case SymbolDeclKind::GlobalLet:
             return out << "global let";
+        case SymbolDeclKind::GlobalConst:
+            return out << "global const";
         case SymbolDeclKind::Alias:
             return out << "alias";
         case SymbolDeclKind::Struct:
@@ -235,6 +250,18 @@ std::ostream& operator<<(std::ostream& out, SymbolUseKind kind) {
             return out << "global let vector element type";
         case SymbolUseKind::GlobalLetMatrixElemType:
             return out << "global let matrix element type";
+        case SymbolUseKind::GlobalConstType:
+            return out << "global const type";
+        case SymbolUseKind::GlobalConstValue:
+            return out << "global const value";
+        case SymbolUseKind::GlobalConstArrayElemType:
+            return out << "global const array element type";
+        case SymbolUseKind::GlobalConstArraySizeValue:
+            return out << "global const array size value";
+        case SymbolUseKind::GlobalConstVectorElemType:
+            return out << "global const vector element type";
+        case SymbolUseKind::GlobalConstMatrixElemType:
+            return out << "global const matrix element type";
         case SymbolUseKind::AliasType:
             return out << "alias type";
         case SymbolUseKind::StructMemberType:
@@ -286,6 +313,10 @@ std::string DiagString(SymbolUseKind kind) {
         case SymbolUseKind::GlobalLetArrayElemType:
         case SymbolUseKind::GlobalLetVectorElemType:
         case SymbolUseKind::GlobalLetMatrixElemType:
+        case SymbolUseKind::GlobalConstType:
+        case SymbolUseKind::GlobalConstArrayElemType:
+        case SymbolUseKind::GlobalConstVectorElemType:
+        case SymbolUseKind::GlobalConstMatrixElemType:
         case SymbolUseKind::AliasType:
         case SymbolUseKind::StructMemberType:
         case SymbolUseKind::ParameterType:
@@ -301,6 +332,8 @@ std::string DiagString(SymbolUseKind kind) {
         case SymbolUseKind::GlobalVarArraySizeValue:
         case SymbolUseKind::GlobalLetValue:
         case SymbolUseKind::GlobalLetArraySizeValue:
+        case SymbolUseKind::GlobalConstValue:
+        case SymbolUseKind::GlobalConstArraySizeValue:
         case SymbolUseKind::LocalVarValue:
         case SymbolUseKind::LocalVarArraySizeValue:
         case SymbolUseKind::LocalLetValue:
@@ -321,6 +354,7 @@ int ScopeDepth(SymbolDeclKind kind) {
     switch (kind) {
         case SymbolDeclKind::GlobalVar:
         case SymbolDeclKind::GlobalLet:
+        case SymbolDeclKind::GlobalConst:
         case SymbolDeclKind::Alias:
         case SymbolDeclKind::Struct:
         case SymbolDeclKind::Function:
@@ -355,6 +389,12 @@ int ScopeDepth(SymbolUseKind kind) {
         case SymbolUseKind::GlobalLetArraySizeValue:
         case SymbolUseKind::GlobalLetVectorElemType:
         case SymbolUseKind::GlobalLetMatrixElemType:
+        case SymbolUseKind::GlobalConstType:
+        case SymbolUseKind::GlobalConstValue:
+        case SymbolUseKind::GlobalConstArrayElemType:
+        case SymbolUseKind::GlobalConstArraySizeValue:
+        case SymbolUseKind::GlobalConstVectorElemType:
+        case SymbolUseKind::GlobalConstMatrixElemType:
         case SymbolUseKind::AliasType:
         case SymbolUseKind::StructMemberType:
         case SymbolUseKind::WorkgroupSizeValue:
@@ -428,6 +468,8 @@ const ast::Node* SymbolTestHelper::Add(SymbolDeclKind kind, Symbol symbol, Sourc
             return b.GlobalVar(source, symbol, b.ty.i32(), ast::StorageClass::kPrivate);
         case SymbolDeclKind::GlobalLet:
             return b.GlobalLet(source, symbol, b.ty.i32(), b.Expr(1_i));
+        case SymbolDeclKind::GlobalConst:
+            return b.GlobalConst(source, symbol, b.ty.i32(), b.Expr(1_i));
         case SymbolDeclKind::Alias:
             return b.Alias(source, symbol, b.ty.i32());
         case SymbolDeclKind::Struct:
@@ -536,6 +578,36 @@ const ast::Node* SymbolTestHelper::Add(SymbolUseKind kind, Symbol symbol, Source
             b.GlobalLet(b.Sym(), b.ty.i32(), node);
             return node;
         }
+        case SymbolUseKind::GlobalConstType: {
+            auto* node = b.ty.type_name(source, symbol);
+            b.GlobalConst(b.Sym(), node, b.Expr(1_i));
+            return node;
+        }
+        case SymbolUseKind::GlobalConstArrayElemType: {
+            auto* node = b.ty.type_name(source, symbol);
+            b.GlobalConst(b.Sym(), b.ty.array(node, 4_i), b.Expr(1_i));
+            return node;
+        }
+        case SymbolUseKind::GlobalConstArraySizeValue: {
+            auto* node = b.Expr(source, symbol);
+            b.GlobalConst(b.Sym(), b.ty.array(b.ty.i32(), node), b.Expr(1_i));
+            return node;
+        }
+        case SymbolUseKind::GlobalConstVectorElemType: {
+            auto* node = b.ty.type_name(source, symbol);
+            b.GlobalConst(b.Sym(), b.ty.vec3(node), b.Expr(1_i));
+            return node;
+        }
+        case SymbolUseKind::GlobalConstMatrixElemType: {
+            auto* node = b.ty.type_name(source, symbol);
+            b.GlobalConst(b.Sym(), b.ty.mat3x4(node), b.Expr(1_i));
+            return node;
+        }
+        case SymbolUseKind::GlobalConstValue: {
+            auto* node = b.Expr(source, symbol);
+            b.GlobalConst(b.Sym(), b.ty.i32(), node);
+            return node;
+        }
         case SymbolUseKind::AliasType: {
             auto* node = b.ty.type_name(source, symbol);
             b.Alias(b.Sym(), node);
@@ -777,6 +849,7 @@ TEST_F(ResolverDependencyGraphDeclSelfUse, GlobalVar) {
 12:34 note: var 'SYMBOL' references var 'SYMBOL' here)");
 }
 
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverDependencyGraphDeclSelfUse, GlobalLet) {
     const Symbol symbol = Sym("SYMBOL");
     GlobalLet(symbol, ty.i32(), Mul(Expr(Source{{12, 34}}, symbol), 123_i));
@@ -784,6 +857,13 @@ TEST_F(ResolverDependencyGraphDeclSelfUse, GlobalLet) {
 12:34 note: let 'SYMBOL' references let 'SYMBOL' here)");
 }
 
+TEST_F(ResolverDependencyGraphDeclSelfUse, GlobalConst) {
+    const Symbol symbol = Sym("SYMBOL");
+    GlobalConst(symbol, ty.i32(), Mul(Expr(Source{{12, 34}}, symbol), 123_i));
+    Build(R"(error: cyclic dependency found: 'SYMBOL' -> 'SYMBOL'
+12:34 note: const 'SYMBOL' references const 'SYMBOL' here)");
+}
+
 TEST_F(ResolverDependencyGraphDeclSelfUse, LocalVar) {
     const Symbol symbol = Sym("SYMBOL");
     WrapInFunction(Decl(Var(symbol, ty.i32(), Mul(Expr(Source{{12, 34}}, symbol), 123_i))));
@@ -907,6 +987,18 @@ TEST_F(ResolverDependencyGraphCyclicRefTest, GlobalVar_Direct) {
 56:78 note: var 'V' references var 'V' here)");
 }
 
+TEST_F(ResolverDependencyGraphCyclicRefTest, GlobalConst_Direct) {
+    // let V : i32 = V;
+
+    GlobalConst(Source{{12, 34}}, "V", ty.i32(), Expr(Source{{56, 78}}, "V"));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              R"(12:34 error: cyclic dependency found: 'V' -> 'V'
+56:78 note: const 'V' references const 'V' here)");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverDependencyGraphCyclicRefTest, GlobalLet_Direct) {
     // let V : i32 = V;
 
@@ -935,6 +1027,24 @@ TEST_F(ResolverDependencyGraphCyclicRefTest, GlobalVar_Indirect) {
 2:10 note: var 'X' references var 'Y' here)");
 }
 
+TEST_F(ResolverDependencyGraphCyclicRefTest, GlobalConst_Indirect) {
+    // 1: const Y : i32 = Z;
+    // 2: const X : i32 = Y;
+    // 3: const Z : i32 = X;
+
+    GlobalConst(Source{{1, 1}}, "Y", ty.i32(), Expr(Source{{1, 10}}, "Z"));
+    GlobalConst(Source{{2, 1}}, "X", ty.i32(), Expr(Source{{2, 10}}, "Y"));
+    GlobalConst(Source{{3, 1}}, "Z", ty.i32(), Expr(Source{{3, 10}}, "X"));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              R"(1:1 error: cyclic dependency found: 'Y' -> 'Z' -> 'X' -> 'Y'
+1:10 note: const 'Y' references const 'Z' here
+3:10 note: const 'Z' references const 'X' here
+2:10 note: const 'X' references const 'Y' here)");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverDependencyGraphCyclicRefTest, GlobalLet_Indirect) {
     // 1: let Y : i32 = Z;
     // 2: let X : i32 = Y;
@@ -958,7 +1068,8 @@ TEST_F(ResolverDependencyGraphCyclicRefTest, Mixed_RecursiveDependencies) {
     // 3: struct S { a : A };
     // 4: var Z = L;
     // 5: type R = A;
-    // 6: let L : S = Z;
+    // 6: const C : S = Z;
+    // 7: let L : S = C;
 
     Func(Source{{1, 1}}, "F", {}, ty.type_name(Source{{1, 5}}, "R"),
          {Return(Expr(Source{{1, 10}}, "Z"))});
@@ -966,16 +1077,18 @@ TEST_F(ResolverDependencyGraphCyclicRefTest, Mixed_RecursiveDependencies) {
     Structure(Source{{3, 1}}, "S", {Member("a", ty.type_name(Source{{3, 10}}, "A"))});
     GlobalVar(Source{{4, 1}}, "Z", nullptr, Expr(Source{{4, 10}}, "L"));
     Alias(Source{{5, 1}}, "R", ty.type_name(Source{{5, 10}}, "A"));
-    GlobalLet(Source{{6, 1}}, "L", ty.type_name(Source{{5, 5}}, "S"), Expr(Source{{5, 10}}, "Z"));
+    GlobalConst(Source{{6, 1}}, "C", ty.type_name(Source{{5, 5}}, "S"), Expr(Source{{5, 10}}, "Z"));
+    GlobalLet(Source{{7, 1}}, "L", ty.type_name(Source{{5, 5}}, "S"), Expr(Source{{5, 10}}, "C"));
 
     EXPECT_FALSE(r()->Resolve());
     EXPECT_EQ(r()->error(),
               R"(2:1 error: cyclic dependency found: 'A' -> 'S' -> 'A'
 2:10 note: alias 'A' references struct 'S' here
 3:10 note: struct 'S' references alias 'A' here
-4:1 error: cyclic dependency found: 'Z' -> 'L' -> 'Z'
+4:1 error: cyclic dependency found: 'Z' -> 'L' -> 'C' -> 'Z'
 4:10 note: var 'Z' references let 'L' here
-5:10 note: let 'L' references var 'Z' here)");
+5:10 note: let 'L' references const 'C' here
+5:10 note: const 'C' references var 'Z' here)");
 }
 
 }  // namespace recursive_tests
@@ -1225,6 +1338,7 @@ TEST_F(ResolverDependencyGraphTraversalTest, SymbolsReached) {
     Alias(Sym(), T);
     Structure(Sym(), {Member(Sym(), T)});
     GlobalVar(Sym(), T, V);
+    GlobalConst(Sym(), T, V);
     GlobalLet(Sym(), T, V);
     Func(Sym(),              //
          {Param(Sym(), T)},  //
@@ -1291,11 +1405,12 @@ TEST_F(ResolverDependencyGraphTraversalTest, SymbolsReached) {
 }
 
 TEST_F(ResolverDependencyGraphTraversalTest, InferredType) {
-    // Check that the nullptr of the var / let type doesn't make things explode
+    // Check that the nullptr of the var / const / let type doesn't make things explode
     GlobalVar("a", nullptr, Expr(1_i));
-    GlobalLet("b", nullptr, Expr(1_i));
+    GlobalConst("b", nullptr, Expr(1_i));
-    WrapInFunction(Var("c", nullptr, Expr(1_i)),  //
+    GlobalLet("c", nullptr, Expr(1_i));
-                   Let("d", nullptr, Expr(1_i)));
+    WrapInFunction(Var("d", nullptr, Expr(1_i)),  //
+                   Let("e", nullptr, Expr(1_i)));
     Build();
 }
 

src/tint/resolver/function_validation_test.cc

@@ -424,11 +424,35 @@ TEST_F(ResolverFunctionValidationTest, FunctionParamsConst) {
 
     EXPECT_FALSE(r()->Resolve());
     EXPECT_EQ(r()->error(),
-              "12:34 error: cannot assign to function parameter\nnote: 'arg' is "
+              "12:34 error: cannot assign to function parameter\nnote: 'arg' is declared here:");
-              "declared here:");
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_ConstU32) {
+    // const x = 4u;
+    // const x = 8u;
+    // @compute @workgroup_size(x, y, 16u)
+    // fn main() {}
+    auto* x = GlobalConst("x", ty.u32(), Expr(4_u));
+    auto* y = GlobalConst("y", ty.u32(), Expr(8_u));
+    auto* func = Func("main", {}, ty.void_(), {},
+                      {Stage(ast::PipelineStage::kCompute), WorkgroupSize("x", "y", 16_u)});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem_func = Sem().Get(func);
+    auto* sem_x = Sem().Get<sem::GlobalVariable>(x);
+    auto* sem_y = Sem().Get<sem::GlobalVariable>(y);
+
+    ASSERT_NE(sem_func, nullptr);
+    ASSERT_NE(sem_x, nullptr);
+    ASSERT_NE(sem_y, nullptr);
+
+    EXPECT_TRUE(sem_func->DirectlyReferencedGlobals().contains(sem_x));
+    EXPECT_TRUE(sem_func->DirectlyReferencedGlobals().contains(sem_y));
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_LetU32) {
     // let x = 4u;
     // let x = 8u;
     // @compute @workgroup_size(x, y, 16u)
@@ -517,6 +541,20 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_MismatchType_I32) {
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_TypeMismatch) {
+    // const x = 64u;
+    // @compute @workgroup_size(1i, x)
+    // fn main() {}
+    GlobalConst("x", ty.u32(), Expr(64_u));
+    Func("main", {}, ty.void_(), {},
+         {Stage(ast::PipelineStage::kCompute), WorkgroupSize(Source{{12, 34}}, 1_i, "x")});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              "12:34 error: workgroup_size arguments must be of the same type, either i32 or u32");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Let_TypeMismatch) {
     // let x = 64u;
     // @compute @workgroup_size(1i, x)
     // fn main() {}
@@ -530,6 +568,22 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_TypeMismatch) {
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_TypeMismatch2) {
+    // const x = 64u;
+    // const y = 32i;
+    // @compute @workgroup_size(x, y)
+    // fn main() {}
+    GlobalConst("x", ty.u32(), Expr(64_u));
+    GlobalConst("y", ty.i32(), Expr(32_i));
+    Func("main", {}, ty.void_(), {},
+         {Stage(ast::PipelineStage::kCompute), WorkgroupSize(Source{{12, 34}}, "x", "y")});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              "12:34 error: workgroup_size arguments must be of the same type, either i32 or u32");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Let_TypeMismatch2) {
     // let x = 64u;
     // let y = 32i;
     // @compute @workgroup_size(x, y)
@@ -543,7 +597,24 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_TypeMismatch2) {
     EXPECT_EQ(r()->error(),
               "12:34 error: workgroup_size arguments must be of the same type, either i32 or u32");
 }
+
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Mismatch_ConstU32) {
+    // const x = 4u;
+    // const x = 8u;
+    // @compute @workgroup_size(x, y, 16i)
+    // fn main() {}
+    GlobalConst("x", ty.u32(), Expr(4_u));
+    GlobalConst("y", ty.u32(), Expr(8_u));
+    Func("main", {}, ty.void_(), {},
+         {Stage(ast::PipelineStage::kCompute), WorkgroupSize(Source{{12, 34}}, "x", "y", 16_i)});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              "12:34 error: workgroup_size arguments must be of the same type, either i32 or u32");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Mismatch_LetU32) {
     // let x = 4u;
     // let x = 8u;
     // @compute @workgroup_size(x, y, 16i)
@@ -594,6 +665,21 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Literal_Zero) {
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_BadType) {
+    // const x = 64.0;
+    // @compute @workgroup_size(x)
+    // fn main() {}
+    GlobalConst("x", ty.f32(), Expr(64_f));
+    Func("main", {}, ty.void_(), {},
+         {Stage(ast::PipelineStage::kCompute), WorkgroupSize(Expr(Source{{12, 34}}, "x"))});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              "12:34 error: workgroup_size argument must be either literal or "
+              "module-scope constant of type i32 or u32");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Let_BadType) {
     // let x = 64.0;
     // @compute @workgroup_size(x)
     // fn main() {}
@@ -608,6 +694,19 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_BadType) {
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_Negative) {
+    // const x = -2i;
+    // @compute @workgroup_size(x)
+    // fn main() {}
+    GlobalConst("x", ty.i32(), Expr(-2_i));
+    Func("main", {}, ty.void_(), {},
+         {Stage(ast::PipelineStage::kCompute), WorkgroupSize(Expr(Source{{12, 34}}, "x"))});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "12:34 error: workgroup_size argument must be at least 1");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Let_Negative) {
     // let x = -2i;
     // @compute @workgroup_size(x)
     // fn main() {}
@@ -620,6 +719,19 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_Negative) {
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_Zero) {
+    // const x = 0i;
+    // @compute @workgroup_size(x)
+    // fn main() {}
+    GlobalConst("x", ty.i32(), Expr(0_i));
+    Func("main", {}, ty.void_(), {},
+         {Stage(ast::PipelineStage::kCompute), WorkgroupSize(Expr(Source{{12, 34}}, "x"))});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "12:34 error: workgroup_size argument must be at least 1");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Let_Zero) {
     // let x = 0i;
     // @compute @workgroup_size(x)
     // fn main() {}
@@ -632,6 +744,19 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_Zero) {
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Const_NestedZeroValueConstructor) {
+    // const x = i32(i32(i32()));
+    // @compute @workgroup_size(x)
+    // fn main() {}
+    GlobalConst("x", ty.i32(), Construct(ty.i32(), Construct(ty.i32(), Construct(ty.i32()))));
+    Func("main", {}, ty.void_(), {},
+         {Stage(ast::PipelineStage::kCompute), WorkgroupSize(Expr(Source{{12, 34}}, "x"))});
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "12:34 error: workgroup_size argument must be at least 1");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverFunctionValidationTest, WorkgroupSize_Let_NestedZeroValueConstructor) {
     // let x = i32(i32(i32()));
     // @compute @workgroup_size(x)
     // fn main() {}

src/tint/resolver/inferred_type_test.cc

@@ -75,6 +75,21 @@ Params all_cases[] = {
 
 using ResolverInferredTypeParamTest = ResolverTestWithParam<Params>;
 
+TEST_P(ResolverInferredTypeParamTest, GlobalConst_Pass) {
+    auto& params = GetParam();
+
+    auto* expected_type = params.create_expected_type(*this);
+
+    // const a = <type constructor>;
+    auto* ctor_expr = params.create_value(*this, 0);
+    auto* a = GlobalConst("a", nullptr, ctor_expr);
+    WrapInFunction();
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(TypeOf(a), expected_type);
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_P(ResolverInferredTypeParamTest, GlobalLet_Pass) {
     auto& params = GetParam();
 

src/tint/resolver/pipeline_overridable_constant_test.cc

@@ -38,7 +38,7 @@ class ResolverPipelineOverridableConstantTest : public ResolverTest {
 };
 
 TEST_F(ResolverPipelineOverridableConstantTest, NonOverridable) {
-    auto* a = GlobalLet("a", ty.f32(), Expr(1_f));
+    auto* a = GlobalConst("a", ty.f32(), Expr(1_f));
 
     EXPECT_TRUE(r()->Resolve()) << r()->error();
 

src/tint/resolver/resolver.cc

@@ -318,6 +318,7 @@ sem::Variable* Resolver::Variable(const ast::Variable* v, bool is_global) {
         [&](const ast::Var* var) { return Var(var, is_global); },
         [&](const ast::Let* let) { return Let(let, is_global); },
         [&](const ast::Override* override) { return Override(override); },
+        [&](const ast::Const* const_) { return Const(const_, is_global); },
         [&](Default) {
             TINT_ICE(Resolver, diagnostics_)
                 << "Resolver::GlobalVariable() called with a unknown variable type: "
@@ -433,6 +434,74 @@ sem::Variable* Resolver::Override(const ast::Override* v) {
     return sem;
 }
 
+sem::Variable* Resolver::Const(const ast::Const* c, bool is_global) {
+    const sem::Type* ty = nullptr;
+
+    // If the variable has a declared type, resolve it.
+    if (c->type) {
+        ty = Type(c->type);
+        if (!ty) {
+            return nullptr;
+        }
+    }
+
+    if (!c->constructor) {
+        AddError("'const' declaration must have an initializer", c->source);
+        return nullptr;
+    }
+
+    const auto* rhs = Expression(c->constructor);
+    if (!rhs) {
+        return nullptr;
+    }
+
+    if (ty) {
+        // If an explicit type was specified, materialize to that type
+        rhs = Materialize(rhs, ty);
+    } else {
+        // If no type was specified, infer it from the RHS
+        ty = rhs->Type();
+    }
+
+    const auto value = rhs->ConstantValue();
+    if (!value) {
+        AddError("'const' initializer must be constant expression", c->constructor->source);
+        return nullptr;
+    }
+
+    // TODO(crbug.com/tint/1580): Temporary seatbelt to used to ensure that a `let` cannot be used
+    // to initialize a 'const'. Once we fully implement `const`, and remove constant evaluation from
+    // 'let', this can be removed.
+    if (auto* user = rhs->UnwrapMaterialize()->As<sem::VariableUser>();
+        user && user->Variable()->Is<sem::LocalVariable>() &&
+        user->Variable()->Declaration()->Is<ast::Let>()) {
+        AddError("'const' initializer must be constant expression", c->constructor->source);
+        return nullptr;
+    }
+
+    if (rhs &&
+        !validator_.VariableConstructorOrCast(c, ast::StorageClass::kNone, ty, rhs->Type())) {
+        return nullptr;
+    }
+
+    if (!ApplyStorageClassUsageToType(ast::StorageClass::kNone, const_cast<sem::Type*>(ty),
+                                      c->source)) {
+        AddNote("while instantiating 'const' " + builder_->Symbols().NameFor(c->symbol), c->source);
+        return nullptr;
+    }
+
+    auto* sem = is_global ? static_cast<sem::Variable*>(builder_->create<sem::GlobalVariable>(
+                                c, ty, ast::StorageClass::kNone, ast::Access::kUndefined, value,
+                                sem::BindingPoint{}))
+                          : static_cast<sem::Variable*>(builder_->create<sem::LocalVariable>(
+                                c, ty, ast::StorageClass::kNone, ast::Access::kUndefined,
+                                current_statement_, value));
+
+    sem->SetConstructor(rhs);
+    builder_->Sem().Add(c, sem);
+    return sem;
+}
+
 sem::Variable* Resolver::Var(const ast::Var* var, bool is_global) {
     const sem::Type* storage_ty = nullptr;
 
@@ -865,7 +934,7 @@ bool Resolver::WorkgroupSize(const ast::Function* func) {
         if (auto* user = args[i]->As<sem::VariableUser>()) {
             // We have an variable of a module-scope constant.
             auto* decl = user->Variable()->Declaration();
-            if (!decl->IsAnyOf<ast::Let, ast::Override>()) {
+            if (!decl->IsAnyOf<ast::Let, ast::Const, ast::Override>()) {
                 AddError(kErrBadType, values[i]->source);
                 return false;
             }

src/tint/resolver/resolver.h

@@ -344,6 +344,14 @@ class Resolver {
     /// @param override the variable
     sem::Variable* Override(const ast::Override* override);
 
+    /// @returns the semantic info for an `ast::Const` `v`. If an error is raised, nullptr is
+    /// returned.
+    /// @note this method does not resolve the attributes as these are context-dependent (global,
+    /// local)
+    /// @param const_ the variable
+    /// @param is_global true if this is module scope, otherwise function scope
+    sem::Variable* Const(const ast::Const* const_, bool is_global);
+
     /// @returns the semantic info for the `ast::Var` `var`. If an error is raised, nullptr is
     /// returned.
     /// @note this method does not resolve the attributes as these are context-dependent (global,

src/tint/resolver/resolver_test.cc

@@ -448,8 +448,24 @@ TEST_F(ResolverTest, ArraySize_SignedLiteral) {
     EXPECT_EQ(ary->Count(), 10u);
 }
 
-TEST_F(ResolverTest, ArraySize_UnsignedConstant) {
+TEST_F(ResolverTest, ArraySize_UnsignedConst) {
-    // let size = 0u;
+    // const size = 10u;
+    // var<private> a : array<f32, size>;
+    GlobalConst("size", nullptr, Expr(10_u));
+    auto* a = GlobalVar("a", ty.array(ty.f32(), Expr("size")), ast::StorageClass::kPrivate);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    ASSERT_NE(TypeOf(a), nullptr);
+    auto* ref = TypeOf(a)->As<sem::Reference>();
+    ASSERT_NE(ref, nullptr);
+    auto* ary = ref->StoreType()->As<sem::Array>();
+    EXPECT_EQ(ary->Count(), 10u);
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverTest, ArraySize_UnsignedLet) {
+    // let size = 10u;
     // var<private> a : array<f32, size>;
     GlobalLet("size", nullptr, Expr(10_u));
     auto* a = GlobalVar("a", ty.array(ty.f32(), Expr("size")), ast::StorageClass::kPrivate);
@@ -463,7 +479,23 @@ TEST_F(ResolverTest, ArraySize_UnsignedConstant) {
     EXPECT_EQ(ary->Count(), 10u);
 }
 
-TEST_F(ResolverTest, ArraySize_SignedConstant) {
+TEST_F(ResolverTest, ArraySize_SignedConst) {
+    // const size = 0;
+    // var<private> a : array<f32, size>;
+    GlobalConst("size", nullptr, Expr(10_i));
+    auto* a = GlobalVar("a", ty.array(ty.f32(), Expr("size")), ast::StorageClass::kPrivate);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    ASSERT_NE(TypeOf(a), nullptr);
+    auto* ref = TypeOf(a)->As<sem::Reference>();
+    ASSERT_NE(ref, nullptr);
+    auto* ary = ref->StoreType()->As<sem::Array>();
+    EXPECT_EQ(ary->Count(), 10u);
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverTest, ArraySize_SignedLet) {
     // let size = 0;
     // var<private> a : array<f32, size>;
     GlobalLet("size", nullptr, Expr(10_i));
@@ -615,7 +647,23 @@ TEST_F(ResolverTest, Expr_Identifier_GlobalVariable) {
     EXPECT_EQ(VarOf(ident)->Declaration(), my_var);
 }
 
-TEST_F(ResolverTest, Expr_Identifier_GlobalConstant) {
+TEST_F(ResolverTest, Expr_Identifier_GlobalConst) {
+    auto* my_var = GlobalConst("my_var", ty.f32(), Construct(ty.f32()));
+
+    auto* ident = Expr("my_var");
+    WrapInFunction(ident);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    ASSERT_NE(TypeOf(ident), nullptr);
+    EXPECT_TRUE(TypeOf(ident)->Is<sem::F32>());
+    EXPECT_TRUE(CheckVarUsers(my_var, {ident}));
+    ASSERT_NE(VarOf(ident), nullptr);
+    EXPECT_EQ(VarOf(ident)->Declaration(), my_var);
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverTest, Expr_Identifier_GlobalLet) {
     auto* my_var = GlobalLet("my_var", ty.f32(), Construct(ty.f32()));
 
     auto* ident = Expr("my_var");
@@ -947,7 +995,36 @@ TEST_F(ResolverTest, Function_WorkgroupSize_Literals) {
     EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
 }
 
-TEST_F(ResolverTest, Function_WorkgroupSize_Consts) {
+TEST_F(ResolverTest, Function_WorkgroupSize_ViaConst) {
+    // const width = 16i;
+    // const height = 8i;
+    // const depth = 2i;
+    // @compute @workgroup_size(width, height, depth)
+    // fn main() {}
+    GlobalConst("width", ty.i32(), Expr(16_i));
+    GlobalConst("height", ty.i32(), Expr(8_i));
+    GlobalConst("depth", ty.i32(), Expr(2_i));
+    auto* func = Func("main", {}, ty.void_(), {},
+                      {
+                          Stage(ast::PipelineStage::kCompute),
+                          WorkgroupSize("width", "height", "depth"),
+                      });
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* func_sem = Sem().Get(func);
+    ASSERT_NE(func_sem, nullptr);
+
+    EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 16u);
+    EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 8u);
+    EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 2u);
+    EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
+    EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, nullptr);
+    EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverTest, Function_WorkgroupSize_ViaLet) {
     // let width = 16i;
     // let height = 8i;
     // let depth = 2i;
@@ -975,7 +1052,36 @@ TEST_F(ResolverTest, Function_WorkgroupSize_Consts) {
     EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
 }
 
-TEST_F(ResolverTest, Function_WorkgroupSize_Consts_NestedInitializer) {
+TEST_F(ResolverTest, Function_WorkgroupSize_ViaConst_NestedInitializer) {
+    // const width = i32(i32(i32(8i)));
+    // const height = i32(i32(i32(4i)));
+    // @compute @workgroup_size(width, height)
+    // fn main() {}
+    GlobalConst("width", ty.i32(),
+                Construct(ty.i32(), Construct(ty.i32(), Construct(ty.i32(), 8_i))));
+    GlobalConst("height", ty.i32(),
+                Construct(ty.i32(), Construct(ty.i32(), Construct(ty.i32(), 4_i))));
+    auto* func = Func("main", {}, ty.void_(), {},
+                      {
+                          Stage(ast::PipelineStage::kCompute),
+                          WorkgroupSize("width", "height"),
+                      });
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* func_sem = Sem().Get(func);
+    ASSERT_NE(func_sem, nullptr);
+
+    EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 8u);
+    EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 4u);
+    EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 1u);
+    EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
+    EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, nullptr);
+    EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverTest, Function_WorkgroupSize_ViaLet_NestedInitializer) {
     // let width = i32(i32(i32(8i)));
     // let height = i32(i32(i32(4i)));
     // @compute @workgroup_size(width, height)
@@ -1060,12 +1166,39 @@ TEST_F(ResolverTest, Function_WorkgroupSize_OverridableConsts_NoInit) {
 }
 
 TEST_F(ResolverTest, Function_WorkgroupSize_Mixed) {
+    // @id(1) override height = 2i;
+    // const depth = 3i;
+    // @compute @workgroup_size(8, height, depth)
+    // fn main() {}
+    auto* height = Override("height", ty.i32(), Expr(2_i), {Id(0)});
+    GlobalConst("depth", ty.i32(), Expr(3_i));
+    auto* func = Func("main", {}, ty.void_(), {},
+                      {
+                          Stage(ast::PipelineStage::kCompute),
+                          WorkgroupSize(8_i, "height", "depth"),
+                      });
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* func_sem = Sem().Get(func);
+    ASSERT_NE(func_sem, nullptr);
+
+    EXPECT_EQ(func_sem->WorkgroupSize()[0].value, 8u);
+    EXPECT_EQ(func_sem->WorkgroupSize()[1].value, 2u);
+    EXPECT_EQ(func_sem->WorkgroupSize()[2].value, 3u);
+    EXPECT_EQ(func_sem->WorkgroupSize()[0].overridable_const, nullptr);
+    EXPECT_EQ(func_sem->WorkgroupSize()[1].overridable_const, height);
+    EXPECT_EQ(func_sem->WorkgroupSize()[2].overridable_const, nullptr);
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverTest, Function_WorkgroupSize_Mixed_Let) {
     // @id(1) override height = 2i;
     // let depth = 3i;
     // @compute @workgroup_size(8, height, depth)
     // fn main() {}
     auto* height = Override("height", ty.i32(), Expr(2_i), {Id(0)});
-    GlobalLet("depth", ty.i32(), Expr(3_i));
+    GlobalConst("depth", ty.i32(), Expr(3_i));
     auto* func = Func("main", {}, ty.void_(), {},
                       {
                           Stage(ast::PipelineStage::kCompute),

src/tint/resolver/source_variable_test.cc

@@ -92,6 +92,18 @@ TEST_F(ResolverSourceVariableTest, GlobalOverride) {
     EXPECT_EQ(Sem().Get(expr)->SourceVariable(), sem_a);
 }
 
+TEST_F(ResolverSourceVariableTest, GlobalConst) {
+    auto* a = GlobalConst("a", ty.f32(), Expr(1_f));
+    auto* expr = Expr(a);
+    WrapInFunction(expr);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem_a = Sem().Get(a);
+    EXPECT_EQ(Sem().Get(expr)->SourceVariable(), sem_a);
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverSourceVariableTest, GlobalLet) {
     auto* a = GlobalLet("a", ty.f32(), Expr(1_f));
     auto* expr = Expr(a);

src/tint/resolver/type_validation_test.cc

@@ -88,8 +88,15 @@ TEST_F(ResolverTypeValidationTest, GlobalVariableWithStorageClass_Pass) {
 }
 
 TEST_F(ResolverTypeValidationTest, GlobalConstNoStorageClass_Pass) {
-    // let global_var: f32;
+    // const global_const: f32 = f32();
-    GlobalLet(Source{{12, 34}}, "global_var", ty.f32(), Construct(ty.f32()));
+    GlobalConst(Source{{12, 34}}, "global_const", ty.f32(), Construct(ty.f32()));
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverTypeValidationTest, GlobalLetNoStorageClass_Pass) {
+    // let global_let: f32;
+    GlobalLet(Source{{12, 34}}, "global_let", ty.f32(), Construct(ty.f32()));
 
     EXPECT_TRUE(r()->Resolve()) << r()->error();
 }
@@ -196,6 +203,15 @@ TEST_F(ResolverTypeValidationTest, ArraySize_SignedLiteral_Pass) {
     EXPECT_TRUE(r()->Resolve()) << r()->error();
 }
 
+TEST_F(ResolverTypeValidationTest, ArraySize_UnsignedConst_Pass) {
+    // const size = 4u;
+    // var<private> a : array<f32, size>;
+    GlobalConst("size", nullptr, Expr(4_u));
+    GlobalVar("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")), ast::StorageClass::kPrivate);
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverTypeValidationTest, ArraySize_UnsignedLet_Pass) {
     // let size = 4u;
     // var<private> a : array<f32, size>;
@@ -204,6 +220,15 @@ TEST_F(ResolverTypeValidationTest, ArraySize_UnsignedLet_Pass) {
     EXPECT_TRUE(r()->Resolve()) << r()->error();
 }
 
+TEST_F(ResolverTypeValidationTest, ArraySize_SignedConst_Pass) {
+    // const size = 4i;
+    // var<private> a : array<f32, size>;
+    GlobalConst("size", nullptr, Expr(4_i));
+    GlobalVar("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")), ast::StorageClass::kPrivate);
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverTypeValidationTest, ArraySize_SignedLet_Pass) {
     // let size = 4i;
     // var<private> a : array<f32, size>;
@@ -240,6 +265,16 @@ TEST_F(ResolverTypeValidationTest, ArraySize_SignedLiteral_Negative) {
     EXPECT_EQ(r()->error(), "12:34 error: array size (-10) must be greater than 0");
 }
 
+TEST_F(ResolverTypeValidationTest, ArraySize_UnsignedConst_Zero) {
+    // const size = 0u;
+    // var<private> a : array<f32, size>;
+    GlobalConst("size", nullptr, Expr(0_u));
+    GlobalVar("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")), ast::StorageClass::kPrivate);
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "12:34 error: array size (0) must be greater than 0");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverTypeValidationTest, ArraySize_UnsignedLet_Zero) {
     // let size = 0u;
     // var<private> a : array<f32, size>;
@@ -249,6 +284,16 @@ TEST_F(ResolverTypeValidationTest, ArraySize_UnsignedLet_Zero) {
     EXPECT_EQ(r()->error(), "12:34 error: array size (0) must be greater than 0");
 }
 
+TEST_F(ResolverTypeValidationTest, ArraySize_SignedConst_Zero) {
+    // const size = 0i;
+    // var<private> a : array<f32, size>;
+    GlobalConst("size", nullptr, Expr(0_i));
+    GlobalVar("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")), ast::StorageClass::kPrivate);
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "12:34 error: array size (0) must be greater than 0");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverTypeValidationTest, ArraySize_SignedLet_Zero) {
     // let size = 0i;
     // var<private> a : array<f32, size>;
@@ -258,6 +303,16 @@ TEST_F(ResolverTypeValidationTest, ArraySize_SignedLet_Zero) {
     EXPECT_EQ(r()->error(), "12:34 error: array size (0) must be greater than 0");
 }
 
+TEST_F(ResolverTypeValidationTest, ArraySize_SignedConst_Negative) {
+    // const size = -10i;
+    // var<private> a : array<f32, size>;
+    GlobalConst("size", nullptr, Expr(-10_i));
+    GlobalVar("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")), ast::StorageClass::kPrivate);
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "12:34 error: array size (-10) must be greater than 0");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverTypeValidationTest, ArraySize_SignedLet_Negative) {
     // let size = -10i;
     // var<private> a : array<f32, size>;
@@ -286,6 +341,18 @@ TEST_F(ResolverTypeValidationTest, ArraySize_IVecLiteral) {
               "'vec2<i32>'");
 }
 
+TEST_F(ResolverTypeValidationTest, ArraySize_FloatConst) {
+    // const size = 10.0;
+    // var<private> a : array<f32, size>;
+    GlobalConst("size", nullptr, Expr(10_f));
+    GlobalVar("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")), ast::StorageClass::kPrivate);
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              "12:34 error: array size must evaluate to a constant integer expression, but is type "
+              "'f32'");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverTypeValidationTest, ArraySize_FloatLet) {
     // let size = 10.0;
     // var<private> a : array<f32, size>;
@@ -297,6 +364,18 @@ TEST_F(ResolverTypeValidationTest, ArraySize_FloatLet) {
               "'f32'");
 }
 
+TEST_F(ResolverTypeValidationTest, ArraySize_IVecConst) {
+    // const size = vec2<i32>(100, 100);
+    // var<private> a : array<f32, size>;
+    GlobalConst("size", nullptr, Construct(ty.vec2<i32>(), 100_i, 100_i));
+    GlobalVar("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")), ast::StorageClass::kPrivate);
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              "12:34 error: array size must evaluate to a constant integer expression, but is type "
+              "'vec2<i32>'");
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverTypeValidationTest, ArraySize_IVecLet) {
     // let size = vec2<i32>(100, 100);
     // var<private> a : array<f32, size>;
@@ -345,6 +424,17 @@ TEST_F(ResolverTypeValidationTest, ArraySize_ModuleVar) {
               "12:34 error: array size must evaluate to a constant integer expression");
 }
 
+TEST_F(ResolverTypeValidationTest, ArraySize_FunctionConst) {
+    // {
+    //   const size = 10;
+    //   var a : array<f32, size>;
+    // }
+    auto* size = Const("size", nullptr, Expr(10_i));
+    auto* a = Var("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")));
+    WrapInFunction(size, a);
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
 TEST_F(ResolverTypeValidationTest, ArraySize_FunctionLet) {
     // {
     //   let size = 10;

src/tint/resolver/uniformity.cc

@@ -982,7 +982,7 @@ class UniformityGraph {
         };
 
         auto name = builder_->Symbols().NameFor(ident->symbol);
-        auto* sem = sem_.Get<sem::VariableUser>(ident)->Variable();
+        auto* sem = sem_.Get(ident)->UnwrapMaterialize()->As<sem::VariableUser>()->Variable();
         auto* node = CreateNode(name + "_ident_expr", ident);
         return Switch(
             sem,

src/tint/resolver/validator.cc

@@ -327,7 +327,13 @@ bool Validator::VariableConstructorOrCast(const ast::Variable* v,
 
     // Value type has to match storage type
     if (storage_ty != value_type) {
-        std::string decl = v->Is<ast::Let>() ? "let" : "var";
+        std::string decl = Switch(
+            v,                                           //
+            [&](const ast::Var*) { return "var"; },      //
+            [&](const ast::Let*) { return "let"; },      //
+            [&](const ast::Const*) { return "const"; },  //
+            [&](Default) { return "<unknown>"; });
+
         AddError("cannot initialize " + decl + " of type '" + sem_.TypeNameOf(storage_ty) +
                      "' with value of type '" + sem_.TypeNameOf(rhs_ty) + "'",
                  v->source);
@@ -528,45 +534,8 @@ bool Validator::GlobalVariable(
     std::unordered_map<uint32_t, const sem::Variable*> constant_ids,
     std::unordered_map<const sem::Type*, const Source&> atomic_composite_info) const {
     auto* decl = global->Declaration();
-    if (!NoDuplicateAttributes(decl->attributes)) {
-        return false;
-    }
-
     bool ok = Switch(
         decl,  //
-        [&](const ast::Override*) {
-            for (auto* attr : decl->attributes) {
-                if (auto* id_attr = attr->As<ast::IdAttribute>()) {
-                    uint32_t id = id_attr->value;
-                    auto it = constant_ids.find(id);
-                    if (it != constant_ids.end() && it->second != global) {
-                        AddError("pipeline constant IDs must be unique", attr->source);
-                        AddNote("a pipeline constant with an ID of " + std::to_string(id) +
-                                    " was previously declared here:",
-                                ast::GetAttribute<ast::IdAttribute>(
-                                    it->second->Declaration()->attributes)
-                                    ->source);
-                        return false;
-                    }
-                    if (id > 65535) {
-                        AddError("pipeline constant IDs must be between 0 and 65535", attr->source);
-                        return false;
-                    }
-                } else {
-                    AddError("attribute is not valid for 'override' declaration", attr->source);
-                    return false;
-                }
-            }
-            return Override(global);
-        },
-        [&](const ast::Let*) {
-            if (!decl->attributes.empty()) {
-                AddError("attribute is not valid for module-scope 'let' declaration",
-                         decl->attributes[0]->source);
-                return false;
-            }
-            return Let(global);
-        },
         [&](const ast::Var* var) {
             if (global->StorageClass() == ast::StorageClass::kNone) {
                 AddError("module-scope 'var' declaration must have a storage class", decl->source);
@@ -602,6 +571,53 @@ bool Validator::GlobalVariable(
             }
 
             return Var(global);
+        },
+        [&](const ast::Let*) {
+            if (!decl->attributes.empty()) {
+                AddError("attribute is not valid for module-scope 'let' declaration",
+                         decl->attributes[0]->source);
+                return false;
+            }
+            return Let(global);
+        },
+        [&](const ast::Override*) {
+            for (auto* attr : decl->attributes) {
+                if (auto* id_attr = attr->As<ast::IdAttribute>()) {
+                    uint32_t id = id_attr->value;
+                    auto it = constant_ids.find(id);
+                    if (it != constant_ids.end() && it->second != global) {
+                        AddError("pipeline constant IDs must be unique", attr->source);
+                        AddNote("a pipeline constant with an ID of " + std::to_string(id) +
+                                    " was previously declared here:",
+                                ast::GetAttribute<ast::IdAttribute>(
+                                    it->second->Declaration()->attributes)
+                                    ->source);
+                        return false;
+                    }
+                    if (id > 65535) {
+                        AddError("pipeline constant IDs must be between 0 and 65535", attr->source);
+                        return false;
+                    }
+                } else {
+                    AddError("attribute is not valid for 'override' declaration", attr->source);
+                    return false;
+                }
+            }
+            return Override(global);
+        },
+        [&](const ast::Const*) {
+            if (!decl->attributes.empty()) {
+                AddError("attribute is not valid for module-scope 'const' declaration",
+                         decl->attributes[0]->source);
+                return false;
+            }
+            return Const(global);
+        },
+        [&](Default) {
+            TINT_ICE(Resolver, diagnostics_)
+                << "Validator::GlobalVariable() called with a unknown variable type: "
+                << decl->TypeInfo().name;
+            return false;
         });
 
     if (!ok) {
@@ -688,6 +704,7 @@ bool Validator::Variable(const sem::Variable* v) const {
         [&](const ast::Var*) { return Var(v); },            //
         [&](const ast::Let*) { return Let(v); },            //
         [&](const ast::Override*) { return Override(v); },  //
+        [&](const ast::Const*) { return true; },            //
         [&](Default) {
             TINT_ICE(Resolver, diagnostics_)
                 << "Validator::Variable() called with a unknown variable type: "
@@ -696,46 +713,6 @@ bool Validator::Variable(const sem::Variable* v) const {
         });
 }
 
-bool Validator::Let(const sem::Variable* v) const {
-    auto* decl = v->Declaration();
-    auto* storage_ty = v->Type()->UnwrapRef();
-
-    if (v->Is<sem::GlobalVariable>()) {
-        auto name = symbols_.NameFor(decl->symbol);
-        if (sem::ParseBuiltinType(name) != sem::BuiltinType::kNone) {
-            AddError("'" + name + "' is a builtin and cannot be redeclared as a 'let'",
-                     decl->source);
-            return false;
-        }
-    }
-
-    if (!(storage_ty->IsConstructible() || storage_ty->Is<sem::Pointer>())) {
-        AddError(sem_.TypeNameOf(storage_ty) + " cannot be used as the type of a 'let'",
-                 decl->source);
-        return false;
-    }
-    return true;
-}
-
-bool Validator::Override(const sem::Variable* v) const {
-    auto* decl = v->Declaration();
-    auto* storage_ty = v->Type()->UnwrapRef();
-
-    auto name = symbols_.NameFor(decl->symbol);
-    if (sem::ParseBuiltinType(name) != sem::BuiltinType::kNone) {
-        AddError("'" + name + "' is a builtin and cannot be redeclared as a 'override'",
-                 decl->source);
-        return false;
-    }
-
-    if (!storage_ty->is_scalar()) {
-        AddError(sem_.TypeNameOf(storage_ty) + " cannot be used as the type of a 'override'",
-                 decl->source);
-        return false;
-    }
-    return true;
-}
-
 bool Validator::Var(const sem::Variable* v) const {
     auto* var = v->Declaration()->As<ast::Var>();
     auto* storage_ty = v->Type()->UnwrapRef();
@@ -783,6 +760,58 @@ bool Validator::Var(const sem::Variable* v) const {
     return true;
 }
 
+bool Validator::Let(const sem::Variable* v) const {
+    auto* decl = v->Declaration();
+    auto* storage_ty = v->Type()->UnwrapRef();
+
+    if (v->Is<sem::GlobalVariable>()) {
+        auto name = symbols_.NameFor(decl->symbol);
+        if (sem::ParseBuiltinType(name) != sem::BuiltinType::kNone) {
+            AddError("'" + name + "' is a builtin and cannot be redeclared as a 'let'",
+                     decl->source);
+            return false;
+        }
+    }
+
+    if (!(storage_ty->IsConstructible() || storage_ty->Is<sem::Pointer>())) {
+        AddError(sem_.TypeNameOf(storage_ty) + " cannot be used as the type of a 'let'",
+                 decl->source);
+        return false;
+    }
+    return true;
+}
+
+bool Validator::Override(const sem::Variable* v) const {
+    auto* decl = v->Declaration();
+    auto* storage_ty = v->Type()->UnwrapRef();
+
+    auto name = symbols_.NameFor(decl->symbol);
+    if (sem::ParseBuiltinType(name) != sem::BuiltinType::kNone) {
+        AddError("'" + name + "' is a builtin and cannot be redeclared as a 'override'",
+                 decl->source);
+        return false;
+    }
+
+    if (!storage_ty->is_scalar()) {
+        AddError(sem_.TypeNameOf(storage_ty) + " cannot be used as the type of a 'override'",
+                 decl->source);
+        return false;
+    }
+    return true;
+}
+
+bool Validator::Const(const sem::Variable* v) const {
+    auto* decl = v->Declaration();
+
+    auto name = symbols_.NameFor(decl->symbol);
+    if (sem::ParseBuiltinType(name) != sem::BuiltinType::kNone) {
+        AddError("'" + name + "' is a builtin and cannot be redeclared as a 'const'", decl->source);
+        return false;
+    }
+
+    return true;
+}
+
 bool Validator::Parameter(const ast::Function* func, const sem::Variable* var) const {
     auto* decl = var->Declaration();
 

src/tint/resolver/validator.h

@@ -358,6 +358,11 @@ class Validator {
     /// @returns true on success, false otherwise.
     bool Variable(const sem::Variable* v) const;
 
+    /// Validates a 'var' variable declaration
+    /// @param v the variable to validate
+    /// @returns true on success, false otherwise.
+    bool Var(const sem::Variable* v) const;
+
     /// Validates a 'let' variable declaration
     /// @param v the variable to validate
     /// @returns true on success, false otherwise.
@@ -368,10 +373,10 @@ class Validator {
     /// @returns true on success, false otherwise.
     bool Override(const sem::Variable* v) const;
 
-    /// Validates a 'var' variable declaration
+    /// Validates a 'const' variable declaration
     /// @param v the variable to validate
     /// @returns true on success, false otherwise.
-    bool Var(const sem::Variable* v) const;
+    bool Const(const sem::Variable* v) const;
 
     /// Validates a variable constructor or cast
     /// @param v the variable to validate

src/tint/resolver/variable_test.cc

@@ -237,8 +237,32 @@ TEST_F(ResolverVariableTest, LocalVar_ShadowsGlobalVar) {
     EXPECT_EQ(user_v->Variable(), global);
 }
 
+TEST_F(ResolverVariableTest, LocalVar_ShadowsGlobalConst) {
+    // const a : i32 = 1i;
+    //
+    // fn X() {
+    //   var a = (a == 123);
+    // }
+
+    auto* g = GlobalConst("a", ty.i32(), Expr(1_i));
+    auto* v = Var("a", nullptr, Expr("a"));
+    Func("F", {}, ty.void_(), {Decl(v)});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* global = Sem().Get(g);
+    auto* local = Sem().Get<sem::LocalVariable>(v);
+    ASSERT_NE(local, nullptr);
+    EXPECT_EQ(local->Shadows(), global);
+
+    auto* user_v = Sem().Get<sem::VariableUser>(local->Declaration()->constructor);
+    ASSERT_NE(user_v, nullptr);
+    EXPECT_EQ(user_v->Variable(), global);
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverVariableTest, LocalVar_ShadowsGlobalLet) {
-    // let a : i32 = 1;
+    // let a : i32 = 1i;
     //
     // fn X() {
     //   var a = (a == 123);
@@ -262,7 +286,7 @@ TEST_F(ResolverVariableTest, LocalVar_ShadowsGlobalLet) {
 
 TEST_F(ResolverVariableTest, LocalVar_ShadowsLocalVar) {
     // fn F() {
-    //   var a : i32; // x
+    //   var a : i32 = 1i; // x
     //   {
     //     var a = a; // y
     //   }
@@ -286,9 +310,35 @@ TEST_F(ResolverVariableTest, LocalVar_ShadowsLocalVar) {
     EXPECT_EQ(user_y->Variable(), local_x);
 }
 
+TEST_F(ResolverVariableTest, LocalVar_ShadowsLocalConst) {
+    // fn F() {
+    //   const a : i32 = 1i;
+    //   {
+    //     var a = (a == 123);
+    //   }
+    // }
+
+    auto* c = Const("a", ty.i32(), Expr(1_i));
+    auto* v = Var("a", nullptr, Expr("a"));
+    Func("X", {}, ty.void_(), {Decl(c), Block(Decl(v))});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* local_c = Sem().Get<sem::LocalVariable>(c);
+    auto* local_v = Sem().Get<sem::LocalVariable>(v);
+
+    ASSERT_NE(local_c, nullptr);
+    ASSERT_NE(local_v, nullptr);
+    EXPECT_EQ(local_v->Shadows(), local_c);
+
+    auto* user_v = Sem().Get<sem::VariableUser>(local_v->Declaration()->constructor);
+    ASSERT_NE(user_v, nullptr);
+    EXPECT_EQ(user_v->Variable(), local_c);
+}
+
 TEST_F(ResolverVariableTest, LocalVar_ShadowsLocalLet) {
     // fn F() {
-    //   let a = 1;
+    //   let a : i32 = 1i;
     //   {
     //     var a = (a == 123);
     //   }
@@ -520,11 +570,35 @@ TEST_F(ResolverVariableTest, LocalLet_ShadowsGlobalVar) {
     EXPECT_EQ(user->Variable(), global);
 }
 
-TEST_F(ResolverVariableTest, LocalLet_ShadowsGlobalLet) {
+TEST_F(ResolverVariableTest, LocalLet_ShadowsGlobalConst) {
-    // let a : i32 = 1;
+    // const a : i32 = 1i;
     //
     // fn F() {
-    //   let a = (a == 321);
+    //   let a = a;
+    // }
+
+    auto* g = GlobalConst("a", ty.i32(), Expr(1_i));
+    auto* l = Let("a", nullptr, Expr("a"));
+    Func("F", {}, ty.void_(), {Decl(l)});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* global = Sem().Get(g);
+    auto* local = Sem().Get<sem::LocalVariable>(l);
+    ASSERT_NE(local, nullptr);
+    EXPECT_EQ(local->Shadows(), global);
+
+    auto* user = Sem().Get<sem::VariableUser>(local->Declaration()->constructor);
+    ASSERT_NE(user, nullptr);
+    EXPECT_EQ(user->Variable(), global);
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverVariableTest, LocalLet_ShadowsGlobalLet) {
+    // let a : i32 = 1i;
+    //
+    // fn F() {
+    //   let a = a;
     // }
 
     auto* g = GlobalLet("a", ty.i32(), Expr(1_i));
@@ -545,7 +619,7 @@ TEST_F(ResolverVariableTest, LocalLet_ShadowsGlobalLet) {
 
 TEST_F(ResolverVariableTest, LocalLet_ShadowsLocalVar) {
     // fn F() {
-    //   var a : i32;
+    //   var a : i32 = 1i;
     //   {
     //     let a = a;
     //   }
@@ -569,11 +643,37 @@ TEST_F(ResolverVariableTest, LocalLet_ShadowsLocalVar) {
     EXPECT_EQ(user->Variable(), local_v);
 }
 
+TEST_F(ResolverVariableTest, LocalLet_ShadowsLocalConst) {
+    // fn X() {
+    //   const a : i32 = 1i; // x
+    //   {
+    //     let a = a; // y
+    //   }
+    // }
+
+    auto* x = Const("a", ty.i32(), Expr(1_i));
+    auto* y = Let("a", nullptr, Expr("a"));
+    Func("X", {}, ty.void_(), {Decl(x), Block(Decl(y))});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* local_x = Sem().Get<sem::LocalVariable>(x);
+    auto* local_y = Sem().Get<sem::LocalVariable>(y);
+
+    ASSERT_NE(local_x, nullptr);
+    ASSERT_NE(local_y, nullptr);
+    EXPECT_EQ(local_y->Shadows(), local_x);
+
+    auto* user = Sem().Get<sem::VariableUser>(local_y->Declaration()->constructor);
+    ASSERT_NE(user, nullptr);
+    EXPECT_EQ(user->Variable(), local_x);
+}
+
 TEST_F(ResolverVariableTest, LocalLet_ShadowsLocalLet) {
     // fn X() {
-    //   let a = 1; // x
+    //   let a : i32 = 1i; // x
     //   {
-    //     let a = (a == 321); // y
+    //     let a = a; // y
     //   }
     // }
 
@@ -620,6 +720,368 @@ TEST_F(ResolverVariableTest, LocalLet_ShadowsParam) {
     EXPECT_EQ(user->Variable(), param);
 }
 
+////////////////////////////////////////////////////////////////////////////////////////////////////
+// Function-scope const
+////////////////////////////////////////////////////////////////////////////////////////////////////
+TEST_F(ResolverVariableTest, LocalConst_ShadowsAlias) {
+    // type a = i32;
+    //
+    // fn F() {
+    //   const a = true;
+    // }
+
+    auto* t = Alias("a", ty.i32());
+    auto* c = Const("a", nullptr, Expr(false));
+    Func("F", {}, ty.void_(), {Decl(c)});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* type_t = Sem().Get(t);
+    auto* local = Sem().Get<sem::LocalVariable>(c);
+    ASSERT_NE(local, nullptr);
+    EXPECT_EQ(local->Shadows(), type_t);
+}
+
+TEST_F(ResolverVariableTest, LocalConst_ShadowsStruct) {
+    // struct a {
+    //   m : i32;
+    // };
+    //
+    // fn F() {
+    //   const a = false;
+    // }
+
+    auto* t = Structure("a", {Member("m", ty.i32())});
+    auto* c = Const("a", nullptr, Expr(false));
+    Func("F", {}, ty.void_(), {Decl(c)});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* type_t = Sem().Get(t);
+    auto* local = Sem().Get<sem::LocalVariable>(c);
+    ASSERT_NE(local, nullptr);
+    EXPECT_EQ(local->Shadows(), type_t);
+}
+
+TEST_F(ResolverVariableTest, LocalConst_ShadowsFunction) {
+    // fn a() {
+    //   const a = false;
+    // }
+
+    auto* c = Const("a", nullptr, Expr(false));
+    auto* fb = Func("a", {}, ty.void_(), {Decl(c)});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* func = Sem().Get(fb);
+    ASSERT_NE(func, nullptr);
+
+    auto* local = Sem().Get<sem::LocalVariable>(c);
+    ASSERT_NE(local, nullptr);
+    EXPECT_EQ(local->Shadows(), func);
+}
+
+TEST_F(ResolverVariableTest, LocalConst_ShadowsGlobalVar) {
+    // var<private> a : i32;
+    //
+    // fn F() {
+    //   const a = 1i;
+    // }
+
+    auto* g = GlobalVar("a", ty.i32(), ast::StorageClass::kPrivate);
+    auto* c = Const("a", nullptr, Expr(1_i));
+    Func("F", {}, ty.void_(), {Decl(c)});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* global = Sem().Get(g);
+    auto* local = Sem().Get<sem::LocalVariable>(c);
+    ASSERT_NE(local, nullptr);
+    EXPECT_EQ(local->Shadows(), global);
+}
+
+TEST_F(ResolverVariableTest, LocalConst_ShadowsGlobalConst) {
+    // const a : i32 = 1i;
+    //
+    // fn F() {
+    //   const a = a;
+    // }
+
+    auto* g = GlobalConst("a", ty.i32(), Expr(1_i));
+    auto* c = Const("a", nullptr, Expr("a"));
+    Func("F", {}, ty.void_(), {Decl(c)});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* global = Sem().Get(g);
+    auto* local = Sem().Get<sem::LocalVariable>(c);
+    ASSERT_NE(local, nullptr);
+    EXPECT_EQ(local->Shadows(), global);
+
+    auto* user = Sem().Get<sem::VariableUser>(local->Declaration()->constructor);
+    ASSERT_NE(user, nullptr);
+    EXPECT_EQ(user->Variable(), global);
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
+TEST_F(ResolverVariableTest, LocalConst_ShadowsGlobalLet) {
+    // let a : i32 = 1i;
+    //
+    // fn F() {
+    //   const a = 1i;
+    // }
+
+    auto* g = GlobalLet("a", ty.i32(), Expr(1_i));
+    auto* c = Const("a", nullptr, Expr("a"));
+    Func("F", {}, ty.void_(), {Decl(c)});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* global = Sem().Get(g);
+    auto* local = Sem().Get<sem::LocalVariable>(c);
+    ASSERT_NE(local, nullptr);
+    EXPECT_EQ(local->Shadows(), global);
+}
+
+TEST_F(ResolverVariableTest, LocalConst_ShadowsLocalVar) {
+    // fn F() {
+    //   var a = 1i;
+    //   {
+    //     const a = 1i;
+    //   }
+    // }
+
+    auto* v = Var("a", ty.i32(), Expr(1_i));
+    auto* c = Const("a", nullptr, Expr(1_i));
+    Func("F", {}, ty.void_(), {Decl(v), Block(Decl(c))});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* local_v = Sem().Get<sem::LocalVariable>(v);
+    auto* local_c = Sem().Get<sem::LocalVariable>(c);
+
+    ASSERT_NE(local_v, nullptr);
+    ASSERT_NE(local_c, nullptr);
+    EXPECT_EQ(local_c->Shadows(), local_v);
+}
+
+TEST_F(ResolverVariableTest, LocalConst_ShadowsLocalConst) {
+    // fn X() {
+    //   const a = 1i; // x
+    //   {
+    //     const a = a; // y
+    //   }
+    // }
+
+    auto* x = Const("a", ty.i32(), Expr(1_i));
+    auto* y = Const("a", nullptr, Expr("a"));
+    Func("X", {}, ty.void_(), {Decl(x), Block(Decl(y))});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* local_x = Sem().Get<sem::LocalVariable>(x);
+    auto* local_y = Sem().Get<sem::LocalVariable>(y);
+
+    ASSERT_NE(local_x, nullptr);
+    ASSERT_NE(local_y, nullptr);
+    EXPECT_EQ(local_y->Shadows(), local_x);
+
+    auto* user = Sem().Get<sem::VariableUser>(local_y->Declaration()->constructor);
+    ASSERT_NE(user, nullptr);
+    EXPECT_EQ(user->Variable(), local_x);
+}
+
+TEST_F(ResolverVariableTest, LocalConst_ShadowsLocalLet) {
+    // fn X() {
+    //   let a = 1i; // x
+    //   {
+    //     const a = 1i; // y
+    //   }
+    // }
+
+    auto* l = Let("a", ty.i32(), Expr(1_i));
+    auto* c = Const("a", nullptr, Expr(1_i));
+    Func("X", {}, ty.void_(), {Decl(l), Block(Decl(c))});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* local_l = Sem().Get<sem::LocalVariable>(l);
+    auto* local_c = Sem().Get<sem::LocalVariable>(c);
+
+    ASSERT_NE(local_l, nullptr);
+    ASSERT_NE(local_c, nullptr);
+    EXPECT_EQ(local_c->Shadows(), local_l);
+}
+
+TEST_F(ResolverVariableTest, LocalConst_ShadowsParam) {
+    // fn F(a : i32) {
+    //   {
+    //     const a = 1i;
+    //   }
+    // }
+
+    auto* p = Param("a", ty.i32());
+    auto* c = Const("a", nullptr, Expr(1_i));
+    Func("X", {p}, ty.void_(), {Block(Decl(c))});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* param = Sem().Get<sem::Parameter>(p);
+    auto* local = Sem().Get<sem::LocalVariable>(c);
+
+    ASSERT_NE(param, nullptr);
+    ASSERT_NE(local, nullptr);
+    EXPECT_EQ(local->Shadows(), param);
+}
+
+TEST_F(ResolverVariableTest, LocalConst_ExplicitType_Decls) {
+    auto* c_i32 = Const("a", ty.i32(), Expr(0_i));
+    auto* c_u32 = Const("b", ty.u32(), Expr(0_u));
+    auto* c_f32 = Const("c", ty.f32(), Expr(0_f));
+    auto* c_vi32 = Const("d", ty.vec3<i32>(), vec3<i32>());
+    auto* c_vu32 = Const("e", ty.vec3<u32>(), vec3<u32>());
+    auto* c_vf32 = Const("f", ty.vec3<f32>(), vec3<f32>());
+    auto* c_mf32 = Const("g", ty.mat3x3<f32>(), mat3x3<f32>());
+
+    WrapInFunction(c_i32, c_u32, c_f32, c_vi32, c_vu32, c_vf32, c_mf32);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    EXPECT_EQ(Sem().Get(c_i32)->Declaration(), c_i32);
+    EXPECT_EQ(Sem().Get(c_u32)->Declaration(), c_u32);
+    EXPECT_EQ(Sem().Get(c_f32)->Declaration(), c_f32);
+    EXPECT_EQ(Sem().Get(c_vi32)->Declaration(), c_vi32);
+    EXPECT_EQ(Sem().Get(c_vu32)->Declaration(), c_vu32);
+    EXPECT_EQ(Sem().Get(c_vf32)->Declaration(), c_vf32);
+    EXPECT_EQ(Sem().Get(c_mf32)->Declaration(), c_mf32);
+
+    ASSERT_TRUE(TypeOf(c_i32)->Is<sem::I32>());
+    ASSERT_TRUE(TypeOf(c_u32)->Is<sem::U32>());
+    ASSERT_TRUE(TypeOf(c_f32)->Is<sem::F32>());
+    ASSERT_TRUE(TypeOf(c_vi32)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_vu32)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_vf32)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_mf32)->Is<sem::Matrix>());
+
+    EXPECT_TRUE(Sem().Get(c_i32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_u32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_f32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vi32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vu32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vf32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_mf32)->ConstantValue().AllZero());
+
+    EXPECT_EQ(Sem().Get(c_i32)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_u32)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_f32)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_vi32)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_vu32)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_vf32)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_mf32)->ConstantValue().ElementCount(), 9u);
+}
+
+TEST_F(ResolverVariableTest, LocalConst_ImplicitType_Decls) {
+    auto* c_i32 = Const("a", nullptr, Expr(0_i));
+    auto* c_u32 = Const("b", nullptr, Expr(0_u));
+    auto* c_f32 = Const("c", nullptr, Expr(0_f));
+    auto* c_ai = Const("d", nullptr, Expr(0_a));
+    auto* c_af = Const("e", nullptr, Expr(0._a));
+    auto* c_vi32 = Const("f", nullptr, vec3<i32>());
+    auto* c_vu32 = Const("g", nullptr, vec3<u32>());
+    auto* c_vf32 = Const("h", nullptr, vec3<f32>());
+    auto* c_vai = Const("i", nullptr, Construct(ty.vec(nullptr, 3), Expr(0_a)));
+    auto* c_vaf = Const("j", nullptr, Construct(ty.vec(nullptr, 3), Expr(0._a)));
+    auto* c_mf32 = Const("k", nullptr, mat3x3<f32>());
+    auto* c_maf32 = Const("l", nullptr, Construct(ty.mat(nullptr, 3, 3), Expr(0._a)));
+
+    WrapInFunction(c_i32, c_u32, c_f32, c_ai, c_af, c_vi32, c_vu32, c_vf32, c_vai, c_vaf, c_mf32,
+                   c_maf32);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    EXPECT_EQ(Sem().Get(c_i32)->Declaration(), c_i32);
+    EXPECT_EQ(Sem().Get(c_u32)->Declaration(), c_u32);
+    EXPECT_EQ(Sem().Get(c_f32)->Declaration(), c_f32);
+    EXPECT_EQ(Sem().Get(c_ai)->Declaration(), c_ai);
+    EXPECT_EQ(Sem().Get(c_af)->Declaration(), c_af);
+    EXPECT_EQ(Sem().Get(c_vi32)->Declaration(), c_vi32);
+    EXPECT_EQ(Sem().Get(c_vu32)->Declaration(), c_vu32);
+    EXPECT_EQ(Sem().Get(c_vf32)->Declaration(), c_vf32);
+    EXPECT_EQ(Sem().Get(c_vai)->Declaration(), c_vai);
+    EXPECT_EQ(Sem().Get(c_vaf)->Declaration(), c_vaf);
+    EXPECT_EQ(Sem().Get(c_mf32)->Declaration(), c_mf32);
+    EXPECT_EQ(Sem().Get(c_maf32)->Declaration(), c_maf32);
+
+    ASSERT_TRUE(TypeOf(c_i32)->Is<sem::I32>());
+    ASSERT_TRUE(TypeOf(c_u32)->Is<sem::U32>());
+    ASSERT_TRUE(TypeOf(c_f32)->Is<sem::F32>());
+    ASSERT_TRUE(TypeOf(c_ai)->Is<sem::AbstractInt>());
+    ASSERT_TRUE(TypeOf(c_af)->Is<sem::AbstractFloat>());
+    ASSERT_TRUE(TypeOf(c_vi32)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_vu32)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_vf32)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_vai)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_vaf)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_mf32)->Is<sem::Matrix>());
+    ASSERT_TRUE(TypeOf(c_maf32)->Is<sem::Matrix>());
+
+    EXPECT_TRUE(Sem().Get(c_i32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_u32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_f32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_ai)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_af)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vi32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vu32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vf32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vai)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vaf)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_mf32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_maf32)->ConstantValue().AllZero());
+
+    EXPECT_EQ(Sem().Get(c_i32)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_u32)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_f32)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_ai)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_af)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_vi32)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_vu32)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_vf32)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_vai)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_vaf)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_mf32)->ConstantValue().ElementCount(), 9u);
+    EXPECT_EQ(Sem().Get(c_maf32)->ConstantValue().ElementCount(), 9u);
+}
+
+// Enable when constants propagate between 'const' variables
+TEST_F(ResolverVariableTest, DISABLED_LocalConst_PropagateConstValue) {
+    auto* a = Const("a", nullptr, Expr(42_i));
+    auto* b = Const("b", nullptr, Expr("a"));
+    auto* c = Const("c", nullptr, Expr("b"));
+
+    WrapInFunction(a, b, c);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    ASSERT_TRUE(TypeOf(c)->Is<sem::I32>());
+
+    ASSERT_EQ(Sem().Get(c)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c)->ConstantValue().Element<i32>(0), 42_i);
+}
+
+// Enable when we have @const operators implemented
+TEST_F(ResolverVariableTest, DISABLED_LocalConst_ConstEval) {
+    auto* c = Const("c", nullptr, Div(Mul(Add(1_i, 2_i), 3_i), 2_i));
+
+    WrapInFunction(c);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    ASSERT_TRUE(TypeOf(c)->Is<sem::I32>());
+
+    ASSERT_EQ(Sem().Get(c)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c)->ConstantValue().Element<i32>(0), 3_i);
+}
+
 ////////////////////////////////////////////////////////////////////////////////////////////////////
 // Module-scope 'var'
 ////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -678,6 +1140,148 @@ TEST_F(ResolverVariableTest, GlobalVar_ExplicitStorageClass) {
     EXPECT_EQ(TypeOf(storage)->As<sem::Reference>()->Access(), ast::Access::kReadWrite);
 }
 
+////////////////////////////////////////////////////////////////////////////////////////////////////
+// Module-scope const
+////////////////////////////////////////////////////////////////////////////////////////////////////
+TEST_F(ResolverVariableTest, GlobalConst_ExplicitType_Decls) {
+    auto* c_i32 = GlobalConst("a", ty.i32(), Expr(0_i));
+    auto* c_u32 = GlobalConst("b", ty.u32(), Expr(0_u));
+    auto* c_f32 = GlobalConst("c", ty.f32(), Expr(0_f));
+    auto* c_vi32 = GlobalConst("d", ty.vec3<i32>(), vec3<i32>());
+    auto* c_vu32 = GlobalConst("e", ty.vec3<u32>(), vec3<u32>());
+    auto* c_vf32 = GlobalConst("f", ty.vec3<f32>(), vec3<f32>());
+    auto* c_mf32 = GlobalConst("g", ty.mat3x3<f32>(), mat3x3<f32>());
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    EXPECT_EQ(Sem().Get(c_i32)->Declaration(), c_i32);
+    EXPECT_EQ(Sem().Get(c_u32)->Declaration(), c_u32);
+    EXPECT_EQ(Sem().Get(c_f32)->Declaration(), c_f32);
+    EXPECT_EQ(Sem().Get(c_vi32)->Declaration(), c_vi32);
+    EXPECT_EQ(Sem().Get(c_vu32)->Declaration(), c_vu32);
+    EXPECT_EQ(Sem().Get(c_vf32)->Declaration(), c_vf32);
+    EXPECT_EQ(Sem().Get(c_mf32)->Declaration(), c_mf32);
+
+    ASSERT_TRUE(TypeOf(c_i32)->Is<sem::I32>());
+    ASSERT_TRUE(TypeOf(c_u32)->Is<sem::U32>());
+    ASSERT_TRUE(TypeOf(c_f32)->Is<sem::F32>());
+    ASSERT_TRUE(TypeOf(c_vi32)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_vu32)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_vf32)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_mf32)->Is<sem::Matrix>());
+
+    EXPECT_TRUE(Sem().Get(c_i32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_u32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_f32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vi32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vu32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vf32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_mf32)->ConstantValue().AllZero());
+
+    EXPECT_EQ(Sem().Get(c_i32)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_u32)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_f32)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_vi32)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_vu32)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_vf32)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_mf32)->ConstantValue().ElementCount(), 9u);
+}
+
+TEST_F(ResolverVariableTest, GlobalConst_ImplicitType_Decls) {
+    auto* c_i32 = GlobalConst("a", nullptr, Expr(0_i));
+    auto* c_u32 = GlobalConst("b", nullptr, Expr(0_u));
+    auto* c_f32 = GlobalConst("c", nullptr, Expr(0_f));
+    auto* c_ai = GlobalConst("d", nullptr, Expr(0_a));
+    auto* c_af = GlobalConst("e", nullptr, Expr(0._a));
+    auto* c_vi32 = GlobalConst("f", nullptr, vec3<i32>());
+    auto* c_vu32 = GlobalConst("g", nullptr, vec3<u32>());
+    auto* c_vf32 = GlobalConst("h", nullptr, vec3<f32>());
+    auto* c_vai = GlobalConst("i", nullptr, Construct(ty.vec(nullptr, 3), Expr(0_a)));
+    auto* c_vaf = GlobalConst("j", nullptr, Construct(ty.vec(nullptr, 3), Expr(0._a)));
+    auto* c_mf32 = GlobalConst("k", nullptr, mat3x3<f32>());
+    auto* c_maf32 = GlobalConst("l", nullptr, Construct(ty.mat(nullptr, 3, 3), Expr(0._a)));
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    EXPECT_EQ(Sem().Get(c_i32)->Declaration(), c_i32);
+    EXPECT_EQ(Sem().Get(c_u32)->Declaration(), c_u32);
+    EXPECT_EQ(Sem().Get(c_f32)->Declaration(), c_f32);
+    EXPECT_EQ(Sem().Get(c_ai)->Declaration(), c_ai);
+    EXPECT_EQ(Sem().Get(c_af)->Declaration(), c_af);
+    EXPECT_EQ(Sem().Get(c_vi32)->Declaration(), c_vi32);
+    EXPECT_EQ(Sem().Get(c_vu32)->Declaration(), c_vu32);
+    EXPECT_EQ(Sem().Get(c_vf32)->Declaration(), c_vf32);
+    EXPECT_EQ(Sem().Get(c_vai)->Declaration(), c_vai);
+    EXPECT_EQ(Sem().Get(c_vaf)->Declaration(), c_vaf);
+    EXPECT_EQ(Sem().Get(c_mf32)->Declaration(), c_mf32);
+    EXPECT_EQ(Sem().Get(c_maf32)->Declaration(), c_maf32);
+
+    ASSERT_TRUE(TypeOf(c_i32)->Is<sem::I32>());
+    ASSERT_TRUE(TypeOf(c_u32)->Is<sem::U32>());
+    ASSERT_TRUE(TypeOf(c_f32)->Is<sem::F32>());
+    ASSERT_TRUE(TypeOf(c_ai)->Is<sem::AbstractInt>());
+    ASSERT_TRUE(TypeOf(c_af)->Is<sem::AbstractFloat>());
+    ASSERT_TRUE(TypeOf(c_vi32)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_vu32)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_vf32)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_vai)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_vaf)->Is<sem::Vector>());
+    ASSERT_TRUE(TypeOf(c_mf32)->Is<sem::Matrix>());
+    ASSERT_TRUE(TypeOf(c_maf32)->Is<sem::Matrix>());
+
+    EXPECT_TRUE(Sem().Get(c_i32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_u32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_f32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_ai)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_af)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vi32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vu32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vf32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vai)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_vaf)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_mf32)->ConstantValue().AllZero());
+    EXPECT_TRUE(Sem().Get(c_maf32)->ConstantValue().AllZero());
+
+    EXPECT_EQ(Sem().Get(c_i32)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_u32)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_f32)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_ai)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_af)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c_vi32)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_vu32)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_vf32)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_vai)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_vaf)->ConstantValue().ElementCount(), 3u);
+    EXPECT_EQ(Sem().Get(c_mf32)->ConstantValue().ElementCount(), 9u);
+    EXPECT_EQ(Sem().Get(c_maf32)->ConstantValue().ElementCount(), 9u);
+}
+
+// Enable when constants propagate between 'const' variables
+TEST_F(ResolverVariableTest, DISABLED_GlobalConst_PropagateConstValue) {
+    GlobalConst("b", nullptr, Expr("a"));
+    auto* c = GlobalConst("c", nullptr, Expr("b"));
+    GlobalConst("a", nullptr, Expr(42_i));
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    ASSERT_TRUE(TypeOf(c)->Is<sem::I32>());
+
+    ASSERT_EQ(Sem().Get(c)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c)->ConstantValue().Element<i32>(0), 42_i);
+}
+
+// Enable when we have @const operators implemented
+TEST_F(ResolverVariableTest, DISABLED_GlobalConst_ConstEval) {
+    auto* c = GlobalConst("c", nullptr, Div(Mul(Add(1_i, 2_i), 3_i), 2_i));
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    ASSERT_TRUE(TypeOf(c)->Is<sem::I32>());
+
+    ASSERT_EQ(Sem().Get(c)->ConstantValue().ElementCount(), 1u);
+    EXPECT_EQ(Sem().Get(c)->ConstantValue().Element<i32>(0), 3_i);
+}
+
 ////////////////////////////////////////////////////////////////////////////////////////////////////
 // Function parameter
 ////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -720,8 +1324,30 @@ TEST_F(ResolverVariableTest, Param_ShadowsGlobalVar) {
     EXPECT_EQ(param->Shadows(), global);
 }
 
+TEST_F(ResolverVariableTest, Param_ShadowsGlobalConst) {
+    // const a : i32 = 1i;
+    //
+    // fn F(a : bool) {
+    // }
+
+    auto* g = GlobalConst("a", ty.i32(), Expr(1_i));
+    auto* p = Param("a", ty.bool_());
+    Func("F", {p}, ty.void_(), {});
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* global = Sem().Get(g);
+    auto* param = Sem().Get<sem::Parameter>(p);
+
+    ASSERT_NE(global, nullptr);
+    ASSERT_NE(param, nullptr);
+
+    EXPECT_EQ(param->Shadows(), global);
+}
+
+// TODO(crbug.com/tint/1580): Remove when module-scope 'let' is removed
 TEST_F(ResolverVariableTest, Param_ShadowsGlobalLet) {
-    // let a : i32 = 1;
+    // let a : i32 = 1i;
     //
     // fn F(a : bool) {
     // }

src/tint/resolver/variable_validation_test.cc

@@ -90,6 +90,15 @@ TEST_F(ResolverVariableValidationTest, OverrideInferedTypeNotScalar) {
     EXPECT_EQ(r()->error(), "56:78 error: vec3<f32> cannot be used as the type of a 'override'");
 }
 
+TEST_F(ResolverVariableValidationTest, ConstConstructorWrongType) {
+    // const c : i32 = 2u
+    WrapInFunction(Const(Source{{3, 3}}, "c", ty.i32(), Expr(2_u)));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              R"(3:3 error: cannot initialize const of type 'i32' with value of type 'u32')");
+}
+
 TEST_F(ResolverVariableValidationTest, LetConstructorWrongType) {
     // var v : i32 = 2u
     WrapInFunction(Let(Source{{3, 3}}, "v", ty.i32(), Expr(2_u)));
@@ -108,6 +117,15 @@ TEST_F(ResolverVariableValidationTest, VarConstructorWrongType) {
               R"(3:3 error: cannot initialize var of type 'i32' with value of type 'u32')");
 }
 
+TEST_F(ResolverVariableValidationTest, ConstConstructorWrongTypeViaAlias) {
+    auto* a = Alias("I32", ty.i32());
+    WrapInFunction(Const(Source{{3, 3}}, "v", ty.Of(a), Expr(2_u)));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(),
+              R"(3:3 error: cannot initialize const of type 'i32' with value of type 'u32')");
+}
+
 TEST_F(ResolverVariableValidationTest, LetConstructorWrongTypeViaAlias) {
     auto* a = Alias("I32", ty.i32());
     WrapInFunction(Let(Source{{3, 3}}, "v", ty.Of(a), Expr(2_u)));
@@ -287,6 +305,14 @@ TEST_F(ResolverVariableValidationTest, InvalidStorageClassForInitializer) {
               "storage classes 'private' and 'function'");
 }
 
+TEST_F(ResolverVariableValidationTest, VectorConstNoType) {
+    // const a : mat3x3 = mat3x3<f32>();
+    WrapInFunction(Const("a", create<ast::Vector>(Source{{12, 34}}, nullptr, 3), vec3<f32>()));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "12:34 error: missing vector element type");
+}
+
 TEST_F(ResolverVariableValidationTest, VectorLetNoType) {
     // let a : mat3x3 = mat3x3<f32>();
     WrapInFunction(Let("a", create<ast::Vector>(Source{{12, 34}}, nullptr, 3), vec3<f32>()));
@@ -303,6 +329,14 @@ TEST_F(ResolverVariableValidationTest, VectorVarNoType) {
     EXPECT_EQ(r()->error(), "12:34 error: missing vector element type");
 }
 
+TEST_F(ResolverVariableValidationTest, MatrixConstNoType) {
+    // const a : mat3x3 = mat3x3<f32>();
+    WrapInFunction(Const("a", create<ast::Matrix>(Source{{12, 34}}, nullptr, 3, 3), mat3x3<f32>()));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), "12:34 error: missing matrix element type");
+}
+
 TEST_F(ResolverVariableValidationTest, MatrixLetNoType) {
     // let a : mat3x3 = mat3x3<f32>();
     WrapInFunction(Let("a", create<ast::Matrix>(Source{{12, 34}}, nullptr, 3, 3), mat3x3<f32>()));
@@ -319,5 +353,49 @@ TEST_F(ResolverVariableValidationTest, MatrixVarNoType) {
     EXPECT_EQ(r()->error(), "12:34 error: missing matrix element type");
 }
 
+TEST_F(ResolverVariableValidationTest, ConstStructure) {
+    auto* s = Structure("S", {Member("m", ty.i32())});
+    auto* c = Const("c", ty.Of(s), Construct(Source{{12, 34}}, ty.Of(s)));
+    WrapInFunction(c);
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), R"(12:34 error: 'const' initializer must be constant expression)");
+}
+
+TEST_F(ResolverVariableValidationTest, GlobalConstStructure) {
+    auto* s = Structure("S", {Member("m", ty.i32())});
+    GlobalConst("c", ty.Of(s), Construct(Source{{12, 34}}, ty.Of(s)));
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), R"(12:34 error: 'const' initializer must be constant expression)");
+}
+
+TEST_F(ResolverVariableValidationTest, ConstInitWithVar) {
+    auto* v = Var("v", nullptr, Expr(1_i));
+    auto* c = Const("c", nullptr, Expr(Source{{12, 34}}, v));
+    WrapInFunction(v, c);
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), R"(12:34 error: 'const' initializer must be constant expression)");
+}
+
+TEST_F(ResolverVariableValidationTest, ConstInitWithOverride) {
+    auto* o = Override("v", nullptr, Expr(1_i));
+    auto* c = Const("c", nullptr, Expr(Source{{12, 34}}, o));
+    WrapInFunction(c);
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), R"(12:34 error: 'const' initializer must be constant expression)");
+}
+
+TEST_F(ResolverVariableValidationTest, ConstInitWithLet) {
+    auto* l = Let("v", nullptr, Expr(1_i));
+    auto* c = Const("c", nullptr, Expr(Source{{12, 34}}, l));
+    WrapInFunction(l, c);
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(r()->error(), R"(12:34 error: 'const' initializer must be constant expression)");
+}
+
 }  // namespace
 }  // namespace tint::resolver