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>
diff --git a/src/tint/program_builder.h b/src/tint/program_builder.h
index 027a0e5..1828391 100644
--- a/src/tint/program_builder.h
+++ b/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"
@@ -1367,6 +1368,35 @@
/// @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
+ /// @param attributes optional variable attributes
/// @returns an `ast::Let` with the given name and type
template <typename NAME>
const ast::Let* Let(NAME&& name,
@@ -1463,6 +1493,42 @@
/// @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
+ /// @param attributes optional variable attributes
/// @returns an `ast::Let` constructed by calling Let() with the arguments of `args`, which is
/// automatically registered as a global variable with the ast::Module.
template <typename NAME>
diff --git a/src/tint/resolver/array_accessor_test.cc b/src/tint/resolver/array_accessor_test.cc
index c0b417c..6aee5e5 100644
--- a/src/tint/resolver/array_accessor_test.cc
+++ b/src/tint/resolver/array_accessor_test.cc
@@ -65,7 +65,7 @@
}
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 @@
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_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) {
- 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, 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();
diff --git a/src/tint/resolver/assignment_validation_test.cc b/src/tint/resolver/assignment_validation_test.cc
index 29e02d5..5c91a76 100644
--- a/src/tint/resolver/assignment_validation_test.cc
+++ b/src/tint/resolver/assignment_validation_test.cc
@@ -61,6 +61,26 @@
}
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, 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>;
diff --git a/src/tint/resolver/attribute_validation_test.cc b/src/tint/resolver/attribute_validation_test.cc
index e593f49..d54ff0e 100644
--- a/src/tint/resolver/attribute_validation_test.cc
+++ b/src/tint/resolver/attribute_validation_test.cc
@@ -774,6 +774,24 @@
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 @@
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,11 +885,11 @@
TestParams{AttributeKind::kBindingAndGroup, false}));
TEST_F(OverrideAttributeTest, DuplicateAttribute) {
- GlobalLet("a", ty.f32(), Expr(1.23_f),
- ast::AttributeList{
- create<ast::IdAttribute>(Source{{12, 34}}, 0),
- create<ast::IdAttribute>(Source{{56, 78}}, 1),
- });
+ Override("a", ty.f32(), Expr(1.23_f),
+ ast::AttributeList{
+ create<ast::IdAttribute>(Source{{12, 34}}, 0),
+ create<ast::IdAttribute>(Source{{56, 78}}, 1),
+ });
WrapInFunction();
diff --git a/src/tint/resolver/builtin_validation_test.cc b/src/tint/resolver/builtin_validation_test.cc
index 61b9e1c..c29aac6 100644
--- a/src/tint/resolver/builtin_validation_test.cc
+++ b/src/tint/resolver/builtin_validation_test.cc
@@ -85,6 +85,15 @@
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, 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);
diff --git a/src/tint/resolver/dependency_graph.cc b/src/tint/resolver/dependency_graph.cc
index 6b3b779..ec450ec 100644
--- a/src/tint/resolver/dependency_graph.cc
+++ b/src/tint/resolver/dependency_graph.cc
@@ -491,9 +491,10 @@
[&](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>";
diff --git a/src/tint/resolver/dependency_graph_test.cc b/src/tint/resolver/dependency_graph_test.cc
index 4ee1dda..ac5117c 100644
--- a/src/tint/resolver/dependency_graph_test.cc
+++ b/src/tint/resolver/dependency_graph_test.cc
@@ -55,6 +55,7 @@
enum class SymbolDeclKind {
GlobalVar,
GlobalLet,
+ GlobalConst,
Alias,
Struct,
Function,
@@ -66,10 +67,10 @@
};
static constexpr SymbolDeclKind kAllSymbolDeclKinds[] = {
- SymbolDeclKind::GlobalVar, SymbolDeclKind::GlobalLet, SymbolDeclKind::Alias,
- SymbolDeclKind::Struct, SymbolDeclKind::Function, SymbolDeclKind::Parameter,
- SymbolDeclKind::LocalVar, SymbolDeclKind::LocalLet, SymbolDeclKind::NestedLocalVar,
- SymbolDeclKind::NestedLocalLet,
+ SymbolDeclKind::GlobalVar, SymbolDeclKind::GlobalLet, SymbolDeclKind::GlobalConst,
+ SymbolDeclKind::Alias, SymbolDeclKind::Struct, SymbolDeclKind::Function,
+ SymbolDeclKind::Parameter, SymbolDeclKind::LocalVar, SymbolDeclKind::LocalLet,
+ SymbolDeclKind::NestedLocalVar, SymbolDeclKind::NestedLocalLet,
};
static constexpr SymbolDeclKind kTypeDeclKinds[] = {
@@ -78,14 +79,14 @@
};
static constexpr SymbolDeclKind kValueDeclKinds[] = {
- SymbolDeclKind::GlobalVar, SymbolDeclKind::GlobalLet, SymbolDeclKind::Parameter,
- SymbolDeclKind::LocalVar, SymbolDeclKind::LocalLet, SymbolDeclKind::NestedLocalVar,
- SymbolDeclKind::NestedLocalLet,
+ SymbolDeclKind::GlobalVar, SymbolDeclKind::GlobalLet, SymbolDeclKind::GlobalConst,
+ SymbolDeclKind::Parameter, SymbolDeclKind::LocalVar, SymbolDeclKind::LocalLet,
+ SymbolDeclKind::NestedLocalVar, SymbolDeclKind::NestedLocalLet,
};
static constexpr SymbolDeclKind kGlobalDeclKinds[] = {
- SymbolDeclKind::GlobalVar, SymbolDeclKind::GlobalLet, SymbolDeclKind::Alias,
- SymbolDeclKind::Struct, SymbolDeclKind::Function,
+ SymbolDeclKind::GlobalVar, SymbolDeclKind::GlobalLet, SymbolDeclKind::GlobalConst,
+ SymbolDeclKind::Alias, SymbolDeclKind::Struct, SymbolDeclKind::Function,
};
static constexpr SymbolDeclKind kLocalDeclKinds[] = {
@@ -96,6 +97,7 @@
static constexpr SymbolDeclKind kGlobalValueDeclKinds[] = {
SymbolDeclKind::GlobalVar,
SymbolDeclKind::GlobalLet,
+ SymbolDeclKind::GlobalConst,
};
static constexpr SymbolDeclKind kFuncDeclKinds[] = {
@@ -119,6 +121,12 @@
GlobalLetVectorElemType,
GlobalLetMatrixElemType,
GlobalLetValue,
+ GlobalConstType,
+ GlobalConstArrayElemType,
+ GlobalConstArraySizeValue,
+ GlobalConstVectorElemType,
+ GlobalConstMatrixElemType,
+ GlobalConstValue,
AliasType,
StructMemberType,
CallFunction,
@@ -151,6 +159,11 @@
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 kValueUseKinds[] = {
SymbolUseKind::GlobalVarValue, SymbolUseKind::GlobalLetValue,
- SymbolUseKind::LocalVarValue, SymbolUseKind::LocalLetValue,
- SymbolUseKind::NestedLocalVarValue, SymbolUseKind::NestedLocalLetValue,
- SymbolUseKind::WorkgroupSizeValue,
+ SymbolUseKind::GlobalConstValue, SymbolUseKind::LocalVarValue,
+ SymbolUseKind::LocalLetValue, SymbolUseKind::NestedLocalVarValue,
+ SymbolUseKind::NestedLocalLetValue, SymbolUseKind::WorkgroupSizeValue,
};
static constexpr SymbolUseKind kFuncUseKinds[] = {
@@ -183,6 +196,8 @@
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 @@
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 @@
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 @@
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 @@
switch (kind) {
case SymbolDeclKind::GlobalVar:
case SymbolDeclKind::GlobalLet:
+ case SymbolDeclKind::GlobalConst:
case SymbolDeclKind::Alias:
case SymbolDeclKind::Struct:
case SymbolDeclKind::Function:
@@ -355,6 +389,12 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
// 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 @@
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 @@
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, 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));
- WrapInFunction(Var("c", nullptr, Expr(1_i)), //
- Let("d", nullptr, Expr(1_i)));
+ GlobalConst("b", nullptr, Expr(1_i));
+ GlobalLet("c", nullptr, Expr(1_i));
+ WrapInFunction(Var("d", nullptr, Expr(1_i)), //
+ Let("e", nullptr, Expr(1_i)));
Build();
}
diff --git a/src/tint/resolver/function_validation_test.cc b/src/tint/resolver/function_validation_test.cc
index 4f801c3..0734dd9 100644
--- a/src/tint/resolver/function_validation_test.cc
+++ b/src/tint/resolver/function_validation_test.cc
@@ -424,11 +424,35 @@
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
- "12:34 error: cannot assign to function parameter\nnote: 'arg' is "
- "declared here:");
+ "12:34 error: cannot assign to function parameter\nnote: 'arg' is 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_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_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 @@
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_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_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_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_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() {}
diff --git a/src/tint/resolver/inferred_type_test.cc b/src/tint/resolver/inferred_type_test.cc
index 1439288..56a22da 100644
--- a/src/tint/resolver/inferred_type_test.cc
+++ b/src/tint/resolver/inferred_type_test.cc
@@ -75,6 +75,21 @@
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();
diff --git a/src/tint/resolver/pipeline_overridable_constant_test.cc b/src/tint/resolver/pipeline_overridable_constant_test.cc
index 9c3b143..583cc16 100644
--- a/src/tint/resolver/pipeline_overridable_constant_test.cc
+++ b/src/tint/resolver/pipeline_overridable_constant_test.cc
@@ -38,7 +38,7 @@
};
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();
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index fc78fb2..b6e5e5b 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -318,6 +318,7 @@
[&](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 @@
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 @@
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;
}
diff --git a/src/tint/resolver/resolver.h b/src/tint/resolver/resolver.h
index 961de41..c162142 100644
--- a/src/tint/resolver/resolver.h
+++ b/src/tint/resolver/resolver.h
@@ -344,6 +344,14 @@
/// @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,
diff --git a/src/tint/resolver/resolver_test.cc b/src/tint/resolver/resolver_test.cc
index d896b90..f4de30c 100644
--- a/src/tint/resolver/resolver_test.cc
+++ b/src/tint/resolver/resolver_test.cc
@@ -448,8 +448,24 @@
EXPECT_EQ(ary->Count(), 10u);
}
-TEST_F(ResolverTest, ArraySize_UnsignedConstant) {
- // let size = 0u;
+TEST_F(ResolverTest, ArraySize_UnsignedConst) {
+ // 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 @@
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 @@
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 @@
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 @@
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)
@@ -1061,11 +1167,38 @@
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),
diff --git a/src/tint/resolver/source_variable_test.cc b/src/tint/resolver/source_variable_test.cc
index 6e6cd43..fbcbba9 100644
--- a/src/tint/resolver/source_variable_test.cc
+++ b/src/tint/resolver/source_variable_test.cc
@@ -92,6 +92,18 @@
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);
diff --git a/src/tint/resolver/type_validation_test.cc b/src/tint/resolver/type_validation_test.cc
index fbf74cf..aaf3b31 100644
--- a/src/tint/resolver/type_validation_test.cc
+++ b/src/tint/resolver/type_validation_test.cc
@@ -88,8 +88,15 @@
}
TEST_F(ResolverTypeValidationTest, GlobalConstNoStorageClass_Pass) {
- // let global_var: f32;
- GlobalLet(Source{{12, 34}}, "global_var", ty.f32(), Construct(ty.f32()));
+ // const global_const: f32 = 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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
"'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 @@
"'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 @@
"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;
diff --git a/src/tint/resolver/uniformity.cc b/src/tint/resolver/uniformity.cc
index 15ca72f..57bf6b0 100644
--- a/src/tint/resolver/uniformity.cc
+++ b/src/tint/resolver/uniformity.cc
@@ -982,7 +982,7 @@
};
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,
diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc
index 75b7f75..5676401 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -327,7 +327,13 @@
// 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 @@
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 @@
}
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 @@
[&](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::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 @@
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();
diff --git a/src/tint/resolver/validator.h b/src/tint/resolver/validator.h
index 9b2a5c6..9610e33 100644
--- a/src/tint/resolver/validator.h
+++ b/src/tint/resolver/validator.h
@@ -358,6 +358,11 @@
/// @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 @@
/// @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
diff --git a/src/tint/resolver/variable_test.cc b/src/tint/resolver/variable_test.cc
index ad3742c..ba62906 100644
--- a/src/tint/resolver/variable_test.cc
+++ b/src/tint/resolver/variable_test.cc
@@ -237,8 +237,32 @@
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_ShadowsLocalVar) {
// fn F() {
- // var a : i32; // x
+ // var a : i32 = 1i; // x
// {
// var a = a; // y
// }
@@ -286,9 +310,35 @@
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 @@
EXPECT_EQ(user->Variable(), global);
}
-TEST_F(ResolverVariableTest, LocalLet_ShadowsGlobalLet) {
- // let a : i32 = 1;
+TEST_F(ResolverVariableTest, LocalLet_ShadowsGlobalConst) {
+ // 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_ShadowsLocalVar) {
// fn F() {
- // var a : i32;
+ // var a : i32 = 1i;
// {
// let a = a;
// }
@@ -569,11 +643,37 @@
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
// }
// }
@@ -621,6 +721,368 @@
}
////////////////////////////////////////////////////////////////////////////////////////////////////
+// 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'
////////////////////////////////////////////////////////////////////////////////////////////////////
TEST_F(ResolverVariableTest, GlobalVar_StorageClass) {
@@ -679,6 +1141,148 @@
}
////////////////////////////////////////////////////////////////////////////////////////////////////
+// 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
////////////////////////////////////////////////////////////////////////////////////////////////////
TEST_F(ResolverVariableTest, Param_ShadowsFunction) {
@@ -720,8 +1324,30 @@
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) {
// }
diff --git a/src/tint/resolver/variable_validation_test.cc b/src/tint/resolver/variable_validation_test.cc
index 60ecaaa..435111b 100644
--- a/src/tint/resolver/variable_validation_test.cc
+++ b/src/tint/resolver/variable_validation_test.cc
@@ -90,6 +90,15 @@
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 @@
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 @@
"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 @@
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 @@
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
