tint/resolver: Track evaluation stage

Keep track of the earliest evaluation point for an expression.
Required to properly track what can be assigned to a `const`, `override`, `let`, `var`.

Bug: tint:1601
Bug: chromium:1343242
Change-Id: I301eec21b71e9036dc1bf6c9af8079317d724762
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/95949
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>

src/tint/BUILD.gn

@@ -420,6 +420,7 @@ libtint_source_set("libtint_core_all_src") {
     "sem/constant.h",
     "sem/depth_multisampled_texture.h",
     "sem/depth_texture.h",
+    "sem/evaluation_stage.h",
     "sem/expression.h",
     "sem/external_texture.h",
     "sem/f16.h",
@@ -621,6 +622,7 @@ libtint_source_set("libtint_sem_src") {
     "sem/depth_multisampled_texture.h",
     "sem/depth_texture.cc",
     "sem/depth_texture.h",
+    "sem/evaluation_stage.h",
     "sem/expression.cc",
     "sem/expression.h",
     "sem/external_texture.cc",
@@ -1093,6 +1095,7 @@ if (tint_build_unittests) {
       "resolver/control_block_validation_test.cc",
       "resolver/dependency_graph_test.cc",
       "resolver/entry_point_validation_test.cc",
+      "resolver/evaluation_stage_test.cc",
       "resolver/function_validation_test.cc",
       "resolver/host_shareable_validation_test.cc",
       "resolver/increment_decrement_validation_test.cc",

src/tint/CMakeLists.txt

@@ -298,6 +298,7 @@ set(TINT_LIB_SRCS
   sem/depth_multisampled_texture.h
   sem/depth_texture.cc
   sem/depth_texture.h
+  sem/evaluation_stage.h
   sem/expression.cc
   sem/expression.h
   sem/external_texture.cc
@@ -777,6 +778,7 @@ if(TINT_BUILD_TESTS)
     resolver/control_block_validation_test.cc
     resolver/dependency_graph_test.cc
     resolver/entry_point_validation_test.cc
+    resolver/evaluation_stage_test.cc
     resolver/function_validation_test.cc
     resolver/host_shareable_validation_test.cc
     resolver/increment_decrement_validation_test.cc

src/tint/resolver/evaluation_stage_test.cc

@@ -0,0 +1,297 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/tint/resolver/resolver.h"
+
+#include "gmock/gmock.h"
+#include "src/tint/resolver/resolver_test_helper.h"
+
+using namespace tint::number_suffixes;  // NOLINT
+
+namespace tint::resolver {
+namespace {
+
+using ResolverEvaluationStageTest = ResolverTest;
+
+TEST_F(ResolverEvaluationStageTest, Literal_i32) {
+    auto* expr = Expr(123_i);
+    WrapInFunction(expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kConstant);
+}
+
+TEST_F(ResolverEvaluationStageTest, Literal_f32) {
+    auto* expr = Expr(123_f);
+    WrapInFunction(expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kConstant);
+}
+
+TEST_F(ResolverEvaluationStageTest, Vector_Ctor) {
+    auto* expr = vec3<f32>();
+    WrapInFunction(expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kConstant);
+}
+
+TEST_F(ResolverEvaluationStageTest, Vector_Ctor_Const_Const) {
+    // const f = 1.f;
+    // vec2<f32>(f, f);
+    auto* f = Const("f", nullptr, Expr(1_f));
+    auto* expr = vec2<f32>(f, f);
+    WrapInFunction(f, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(f)->Stage(), sem::EvaluationStage::kConstant);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kConstant);
+}
+
+TEST_F(ResolverEvaluationStageTest, Vector_Ctor_Runtime_Runtime) {
+    // var f = 1.f;
+    // vec2<f32>(f, f);
+    auto* f = Var("f", nullptr, Expr(1_f));
+    auto* expr = vec2<f32>(f, f);
+    WrapInFunction(f, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(f)->Stage(), sem::EvaluationStage::kRuntime);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kRuntime);
+}
+
+TEST_F(ResolverEvaluationStageTest, Vector_Conv_Const) {
+    // const f = 1.f;
+    // vec2<u32>(vec2<f32>(f));
+    auto* f = Const("f", nullptr, Expr(1_f));
+    auto* expr = vec2<u32>(vec2<f32>(f));
+    WrapInFunction(f, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(f)->Stage(), sem::EvaluationStage::kConstant);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kConstant);
+}
+
+TEST_F(ResolverEvaluationStageTest, Vector_Conv_Runtime) {
+    // var f = 1.f;
+    // vec2<u32>(vec2<f32>(f));
+    auto* f = Var("f", nullptr, Expr(1_f));
+    auto* expr = vec2<u32>(vec2<f32>(f));
+    WrapInFunction(f, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(f)->Stage(), sem::EvaluationStage::kRuntime);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kRuntime);
+}
+
+TEST_F(ResolverEvaluationStageTest, Matrix_Ctor) {
+    auto* expr = mat2x2<f32>();
+    WrapInFunction(expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kConstant);
+}
+
+TEST_F(ResolverEvaluationStageTest, Array_Ctor) {
+    auto* expr = array<f32, 3>();
+    WrapInFunction(expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kConstant);
+}
+
+TEST_F(ResolverEvaluationStageTest, Array_Ctor_Const_Const) {
+    // const f = 1.f;
+    // array<f32, 2>(f, f);
+    auto* f = Const("f", nullptr, Expr(1_f));
+    auto* expr = Construct(ty.array<f32, 2>(), f, f);
+    WrapInFunction(f, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(f)->Stage(), sem::EvaluationStage::kConstant);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kConstant);
+}
+
+TEST_F(ResolverEvaluationStageTest, Array_Ctor_Const_Override) {
+    // const f1 = 1.f;
+    // override f2 = 2.f;
+    // array<f32, 2>(f1, f2);
+    auto* f1 = Const("f1", nullptr, Expr(1_f));
+    auto* f2 = Override("f2", nullptr, Expr(2_f));
+    auto* expr = Construct(ty.array<f32, 2>(), f1, f2);
+    WrapInFunction(f1, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(f1)->Stage(), sem::EvaluationStage::kConstant);
+    EXPECT_EQ(Sem().Get(f2)->Stage(), sem::EvaluationStage::kOverride);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kOverride);
+}
+
+TEST_F(ResolverEvaluationStageTest, Array_Ctor_Override_Runtime) {
+    // override f1 = 1.f;
+    // var f2 = 2.f;
+    // array<f32, 2>(f1, f2);
+    auto* f1 = Override("f1", nullptr, Expr(1_f));
+    auto* f2 = Var("f2", nullptr, Expr(2_f));
+    auto* expr = Construct(ty.array<f32, 2>(), f1, f2);
+    WrapInFunction(f2, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(f1)->Stage(), sem::EvaluationStage::kOverride);
+    EXPECT_EQ(Sem().Get(f2)->Stage(), sem::EvaluationStage::kRuntime);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kRuntime);
+}
+
+TEST_F(ResolverEvaluationStageTest, Array_Ctor_Const_Runtime) {
+    // const f1 = 1.f;
+    // var f2 = 2.f;
+    // array<f32, 2>(f1, f2);
+    auto* f1 = Const("f1", nullptr, Expr(1_f));
+    auto* f2 = Var("f2", nullptr, Expr(2_f));
+    auto* expr = Construct(ty.array<f32, 2>(), f1, f2);
+    WrapInFunction(f1, f2, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(f1)->Stage(), sem::EvaluationStage::kConstant);
+    EXPECT_EQ(Sem().Get(f2)->Stage(), sem::EvaluationStage::kRuntime);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kRuntime);
+}
+
+TEST_F(ResolverEvaluationStageTest, Array_Ctor_Runtime_Runtime) {
+    // var f = 1.f;
+    // array<f32, 2>(f, f);
+    auto* f = Var("f", nullptr, Expr(1_f));
+    auto* expr = Construct(ty.array<f32, 2>(), f, f);
+    WrapInFunction(f, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(f)->Stage(), sem::EvaluationStage::kRuntime);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kRuntime);
+}
+
+TEST_F(ResolverEvaluationStageTest, IndexAccessor_Const_Const) {
+    // const vec = vec4<f32>();
+    // const idx = 1_i;
+    // vec[idx]
+    auto* vec = Const("vec", nullptr, vec4<f32>());
+    auto* idx = Const("idx", nullptr, Expr(1_i));
+    auto* expr = IndexAccessor(vec, idx);
+    WrapInFunction(vec, idx, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(vec)->Stage(), sem::EvaluationStage::kConstant);
+    EXPECT_EQ(Sem().Get(idx)->Stage(), sem::EvaluationStage::kConstant);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kConstant);
+}
+
+TEST_F(ResolverEvaluationStageTest, IndexAccessor_Runtime_Const) {
+    // var vec = vec4<f32>();
+    // const idx = 1_i;
+    // vec[idx]
+    auto* vec = Var("vec", nullptr, vec4<f32>());
+    auto* idx = Const("idx", nullptr, Expr(1_i));
+    auto* expr = IndexAccessor(vec, idx);
+    WrapInFunction(vec, idx, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(vec)->Stage(), sem::EvaluationStage::kRuntime);
+    EXPECT_EQ(Sem().Get(idx)->Stage(), sem::EvaluationStage::kConstant);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kRuntime);
+}
+
+TEST_F(ResolverEvaluationStageTest, IndexAccessor_Const_Override) {
+    // const vec = vec4<f32>();
+    // override idx = 1_i;
+    // vec[idx]
+    auto* vec = Const("vec", nullptr, vec4<f32>());
+    auto* idx = Override("idx", nullptr, Expr(1_i));
+    auto* expr = IndexAccessor(vec, idx);
+    WrapInFunction(vec, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(vec)->Stage(), sem::EvaluationStage::kConstant);
+    EXPECT_EQ(Sem().Get(idx)->Stage(), sem::EvaluationStage::kOverride);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kOverride);
+}
+
+TEST_F(ResolverEvaluationStageTest, IndexAccessor_Const_Runtime) {
+    // const vec = vec4<f32>();
+    // let idx = 1_i;
+    // vec[idx]
+    auto* vec = Const("vec", nullptr, vec4<f32>());
+    auto* idx = Let("idx", nullptr, Expr(1_i));
+    auto* expr = IndexAccessor(vec, idx);
+    WrapInFunction(vec, idx, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(vec)->Stage(), sem::EvaluationStage::kConstant);
+    EXPECT_EQ(Sem().Get(idx)->Stage(), sem::EvaluationStage::kRuntime);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kRuntime);
+}
+
+TEST_F(ResolverEvaluationStageTest, Swizzle_Const) {
+    // const vec = S();
+    // vec.m
+    auto* vec = Const("vec", nullptr, vec4<f32>());
+    auto* expr = MemberAccessor(vec, "xz");
+    WrapInFunction(vec, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(vec)->Stage(), sem::EvaluationStage::kConstant);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kConstant);
+}
+
+TEST_F(ResolverEvaluationStageTest, Swizzle_Runtime) {
+    // var vec = S();
+    // vec.m
+    auto* vec = Var("vec", nullptr, vec4<f32>());
+    auto* expr = MemberAccessor(vec, "rg");
+    WrapInFunction(vec, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(vec)->Stage(), sem::EvaluationStage::kRuntime);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kRuntime);
+}
+
+TEST_F(ResolverEvaluationStageTest, MemberAccessor_Const) {
+    // struct S { m : i32 };
+    // const str = S();
+    // str.m
+    Structure("S", {Member("m", ty.i32())});
+    auto* str = Const("str", nullptr, Construct(ty.type_name("S")));
+    auto* expr = MemberAccessor(str, "m");
+    WrapInFunction(str, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(str)->Stage(), sem::EvaluationStage::kConstant);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kConstant);
+}
+
+TEST_F(ResolverEvaluationStageTest, MemberAccessor_Runtime) {
+    // struct S { m : i32 };
+    // var str = S();
+    // str.m
+    Structure("S", {Member("m", ty.i32())});
+    auto* str = Var("str", nullptr, Construct(ty.type_name("S")));
+    auto* expr = MemberAccessor(str, "m");
+    WrapInFunction(str, expr);
+
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
+    EXPECT_EQ(Sem().Get(str)->Stage(), sem::EvaluationStage::kRuntime);
+    EXPECT_EQ(Sem().Get(expr)->Stage(), sem::EvaluationStage::kRuntime);
+}
+
+}  // namespace
+}  // namespace tint::resolver

src/tint/resolver/intrinsic_table.cc

@@ -1119,8 +1119,10 @@ Impl::Builtin Impl::Lookup(sem::BuiltinType builtin_type,
         if (match.overload->flags.Contains(OverloadFlag::kSupportsComputePipeline)) {
             supported_stages.Add(ast::PipelineStage::kCompute);
         }
+        auto eval_stage = match.overload->const_eval_fn ? sem::EvaluationStage::kConstant
+                                                        : sem::EvaluationStage::kRuntime;
         return builder.create<sem::Builtin>(
-            builtin_type, match.return_type, std::move(params), supported_stages,
+            builtin_type, match.return_type, std::move(params), eval_stage, supported_stages,
             match.overload->flags.Contains(OverloadFlag::kIsDeprecated));
     });
     return Builtin{sem, match.overload->const_eval_fn};
@@ -1292,8 +1294,11 @@ IntrinsicTable::CtorOrConv Impl::Lookup(CtorConvIntrinsic type,
                 nullptr, static_cast<uint32_t>(params.size()), p.type, ast::StorageClass::kNone,
                 ast::Access::kUndefined, p.usage));
         }
+        auto eval_stage = match.overload->const_eval_fn ? sem::EvaluationStage::kConstant
+                                                        : sem::EvaluationStage::kRuntime;
         auto* target = utils::GetOrCreate(constructors, match, [&]() {
-            return builder.create<sem::TypeConstructor>(match.return_type, std::move(params));
+            return builder.create<sem::TypeConstructor>(match.return_type, std::move(params),
+                                                        eval_stage);
         });
         return CtorOrConv{target, match.overload->const_eval_fn};
     }
@@ -1303,7 +1308,9 @@ IntrinsicTable::CtorOrConv Impl::Lookup(CtorConvIntrinsic type,
         auto param = builder.create<sem::Parameter>(
             nullptr, 0u, match.parameters[0].type, ast::StorageClass::kNone,
             ast::Access::kUndefined, match.parameters[0].usage);
-        return builder.create<sem::TypeConversion>(match.return_type, param);
+        auto eval_stage = match.overload->const_eval_fn ? sem::EvaluationStage::kConstant
+                                                        : sem::EvaluationStage::kRuntime;
+        return builder.create<sem::TypeConversion>(match.return_type, param, eval_stage);
     });
     return CtorOrConv{target, match.overload->const_eval_fn};
 }

src/tint/resolver/resolver.cc

@@ -367,10 +367,11 @@ sem::Variable* Resolver::Let(const ast::Let* v, bool is_global) {
     sem::Variable* sem = nullptr;
     if (is_global) {
         sem = builder_->create<sem::GlobalVariable>(
-            v, ty, ast::StorageClass::kNone, ast::Access::kUndefined, /* constant_value */ nullptr,
+            v, ty, sem::EvaluationStage::kRuntime, ast::StorageClass::kNone,
-            sem::BindingPoint{});
+            ast::Access::kUndefined, /* constant_value */ nullptr, sem::BindingPoint{});
     } else {
-        sem = builder_->create<sem::LocalVariable>(v, ty, ast::StorageClass::kNone,
+        sem = builder_->create<sem::LocalVariable>(v, ty, sem::EvaluationStage::kRuntime,
+                                                   ast::StorageClass::kNone,
                                                    ast::Access::kUndefined, current_statement_,
                                                    /* constant_value */ nullptr);
     }
@@ -421,8 +422,8 @@ sem::Variable* Resolver::Override(const ast::Override* v) {
     }
 
     auto* sem = builder_->create<sem::GlobalVariable>(
-        v, ty, ast::StorageClass::kNone, ast::Access::kUndefined, /* constant_value */ nullptr,
+        v, ty, sem::EvaluationStage::kOverride, ast::StorageClass::kNone, ast::Access::kUndefined,
-        sem::BindingPoint{});
+        /* constant_value */ nullptr, sem::BindingPoint{});
 
     if (auto* id = ast::GetAttribute<ast::IdAttribute>(v->attributes)) {
         sem->SetConstantId(static_cast<uint16_t>(id->value));
@@ -482,11 +483,11 @@ sem::Variable* Resolver::Const(const ast::Const* c, bool is_global) {
     }
 
     auto* sem = is_global ? static_cast<sem::Variable*>(builder_->create<sem::GlobalVariable>(
-                                c, ty, ast::StorageClass::kNone, ast::Access::kUndefined, value,
+                                c, ty, sem::EvaluationStage::kConstant, ast::StorageClass::kNone,
-                                sem::BindingPoint{}))
+                                ast::Access::kUndefined, value, sem::BindingPoint{}))
                           : static_cast<sem::Variable*>(builder_->create<sem::LocalVariable>(
-                                c, ty, ast::StorageClass::kNone, ast::Access::kUndefined,
+                                c, ty, sem::EvaluationStage::kConstant, ast::StorageClass::kNone,
-                                current_statement_, value));
+                                ast::Access::kUndefined, current_statement_, value));
 
     sem->SetConstructor(rhs);
     builder_->Sem().Add(c, sem);
@@ -567,12 +568,14 @@ sem::Variable* Resolver::Var(const ast::Var* var, bool is_global) {
         if (auto bp = var->BindingPoint()) {
             binding_point = {bp.group->value, bp.binding->value};
         }
-        sem = builder_->create<sem::GlobalVariable>(var, var_ty, storage_class, access,
+        sem = builder_->create<sem::GlobalVariable>(var, var_ty, sem::EvaluationStage::kRuntime,
+                                                    storage_class, access,
                                                     /* constant_value */ nullptr, binding_point);
 
     } else {
-        sem = builder_->create<sem::LocalVariable>(
+        sem = builder_->create<sem::LocalVariable>(var, var_ty, sem::EvaluationStage::kRuntime,
-            var, var_ty, storage_class, access, current_statement_, /* constant_value */ nullptr);
+                                                   storage_class, access, current_statement_,
+                                                   /* constant_value */ nullptr);
     }
 
     sem->SetConstructor(rhs);
@@ -1270,6 +1273,7 @@ sem::Expression* Resolver::Expression(const ast::Expression* root) {
             [&](const ast::UnaryOpExpression* unary) -> sem::Expression* { return UnaryOp(unary); },
             [&](const ast::PhonyExpression*) -> sem::Expression* {
                 return builder_->create<sem::Expression>(expr, builder_->create<sem::Void>(),
+                                                         sem::EvaluationStage::kRuntime,
                                                          current_statement_,
                                                          /* constant_value */ nullptr,
                                                          /* has_side_effects */ false);
@@ -1423,10 +1427,11 @@ sem::Expression* Resolver::IndexAccessor(const ast::IndexAccessorExpression* exp
         ty = builder_->create<sem::Reference>(ty, ref->StorageClass(), ref->Access());
     }
 
+    auto stage = sem::EarliestStage(obj->Stage(), idx->Stage());
     auto val = const_eval_.Index(obj, idx);
     bool has_side_effects = idx->HasSideEffects() || obj->HasSideEffects();
     auto* sem = builder_->create<sem::IndexAccessorExpression>(
-        expr, ty, obj, idx, current_statement_, std::move(val), has_side_effects,
+        expr, ty, stage, obj, idx, current_statement_, std::move(val), has_side_effects,
         obj->SourceVariable());
     sem->Behaviors() = idx->Behaviors() + obj->Behaviors();
     return sem;
@@ -1443,8 +1448,9 @@ sem::Expression* Resolver::Bitcast(const ast::BitcastExpression* expr) {
     }
 
     auto val = const_eval_.Bitcast(ty, inner);
-    auto* sem = builder_->create<sem::Expression>(expr, ty, current_statement_, std::move(val),
+    auto stage = sem::EvaluationStage::kRuntime;  // TODO(crbug.com/tint/1581)
-                                                  inner->HasSideEffects());
+    auto* sem = builder_->create<sem::Expression>(expr, ty, stage, current_statement_,
+                                                  std::move(val), inner->HasSideEffects());
 
     sem->Behaviors() = inner->Behaviors();
 
@@ -1464,6 +1470,7 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
 
     // Resolve all of the arguments, their types and the set of behaviors.
     std::vector<const sem::Expression*> args(expr->args.size());
+    auto args_stage = sem::EvaluationStage::kConstant;
     sem::Behaviors arg_behaviors;
     for (size_t i = 0; i < expr->args.size(); i++) {
         auto* arg = sem_.Get(expr->args[i]);
@@ -1471,6 +1478,7 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
             return nullptr;
         }
         args[i] = arg;
+        args_stage = sem::EarliestStage(args_stage, arg->Stage());
         arg_behaviors.Add(arg->Behaviors());
     }
     arg_behaviors.Remove(sem::Behavior::kNext);
@@ -1491,14 +1499,38 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
             return nullptr;
         }
         const sem::Constant* value = nullptr;
-        if (ctor_or_conv.const_eval_fn) {
+        auto stage = sem::EarliestStage(ctor_or_conv.target->Stage(), args_stage);
+        if (stage == sem::EvaluationStage::kConstant) {
             value = (const_eval_.*ctor_or_conv.const_eval_fn)(ctor_or_conv.target->ReturnType(),
                                                               args.data(), args.size());
         }
-        return builder_->create<sem::Call>(expr, ctor_or_conv.target, std::move(args),
+        return builder_->create<sem::Call>(expr, ctor_or_conv.target, stage, std::move(args),
                                            current_statement_, value, has_side_effects);
     };
 
+    // ct_ctor_or_conv is a helper for building a sem::TypeConstructor for an array or structure
+    // constructor call target.
+    auto arr_or_str_ctor = [&](const sem::Type* ty,
+                               const sem::CallTarget* call_target) -> sem::Call* {
+        if (!MaterializeArguments(args, call_target)) {
+            return nullptr;
+        }
+
+        auto stage = args_stage;               // The evaluation stage of the call
+        const sem::Constant* value = nullptr;  // The constant value for the call
+        if (stage == sem::EvaluationStage::kConstant) {
+            value = const_eval_.ArrayOrStructCtor(ty, args);
+            if (!value) {
+                // Constant evaluation failed.
+                // Can happen for expressions that will fail validation (later).
+                stage = sem::EvaluationStage::kRuntime;
+            }
+        }
+
+        return builder_->create<sem::Call>(expr, call_target, stage, std::move(args),
+                                           current_statement_, std::move(value), has_side_effects);
+    };
+
     // ct_ctor_or_conv is a helper for building either a sem::TypeConstructor or sem::TypeConversion
     // call for the given semantic type.
     auto ty_ctor_or_conv = [&](const sem::Type* ty) {
@@ -1517,7 +1549,7 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
             [&](const sem::Bool*) { return ct_ctor_or_conv(CtorConvIntrinsic::kBool, nullptr); },
             [&](const sem::Array* arr) -> sem::Call* {
                 auto* call_target = utils::GetOrCreate(
-                    array_ctors_, ArrayConstructorSig{{arr, args.size()}},
+                    array_ctors_, ArrayConstructorSig{{arr, args.size(), args_stage}},
                     [&]() -> sem::TypeConstructor* {
                         sem::ParameterList params(args.size());
                         for (size_t i = 0; i < args.size(); i++) {
@@ -1528,18 +1560,14 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
                                 ast::StorageClass::kNone,  // storage_class
                                 ast::Access::kUndefined);  // access
                         }
-                        return builder_->create<sem::TypeConstructor>(arr, std::move(params));
+                        return builder_->create<sem::TypeConstructor>(arr, std::move(params),
+                                                                      args_stage);
                     });
-                if (!MaterializeArguments(args, call_target)) {
+                return arr_or_str_ctor(arr, call_target);
-                    return nullptr;
-                }
-                auto val = const_eval_.ArrayOrStructCtor(arr, args);
-                return builder_->create<sem::Call>(expr, call_target, std::move(args),
-                                                   current_statement_, val, has_side_effects);
             },
             [&](const sem::Struct* str) -> sem::Call* {
                 auto* call_target = utils::GetOrCreate(
-                    struct_ctors_, StructConstructorSig{{str, args.size()}},
+                    struct_ctors_, StructConstructorSig{{str, args.size(), args_stage}},
                     [&]() -> sem::TypeConstructor* {
                         sem::ParameterList params(std::min(args.size(), str->Members().size()));
                         for (size_t i = 0, n = params.size(); i < n; i++) {
@@ -1550,15 +1578,10 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
                                 ast::StorageClass::kNone,   // storage_class
                                 ast::Access::kUndefined);   // access
                         }
-                        return builder_->create<sem::TypeConstructor>(str, std::move(params));
+                        return builder_->create<sem::TypeConstructor>(str, std::move(params),
+                                                                      args_stage);
                     });
-                if (!MaterializeArguments(args, call_target)) {
+                return arr_or_str_ctor(str, call_target);
-                    return nullptr;
-                }
-                auto val = const_eval_.ArrayOrStructCtor(str, args);
-                return builder_->create<sem::Call>(expr, call_target, std::move(args),
-                                                   current_statement_, std::move(val),
-                                                   has_side_effects);
             },
             [&](Default) {
                 AddError("type is not constructible", expr->source);
@@ -1686,9 +1709,16 @@ sem::Call* Resolver::BuiltinCall(const ast::CallExpression* expr,
         AddWarning("use of deprecated builtin", expr->source);
     }
 
+    auto stage = builtin.sem->Stage();
+    if (stage == sem::EvaluationStage::kConstant) {
+        for (auto* arg : args) {
+            stage = sem::EarliestStage(stage, arg->Stage());
+        }
+    }
+
     // If the builtin is @const, and all arguments have constant values, evaluate the builtin now.
     const sem::Constant* value = nullptr;
-    if (builtin.const_eval_fn) {
+    if (stage == sem::EvaluationStage::kConstant && builtin.const_eval_fn) {
         value = (const_eval_.*builtin.const_eval_fn)(builtin.sem->ReturnType(), args.data(),
                                                      args.size());
     }
@@ -1696,8 +1726,8 @@ sem::Call* Resolver::BuiltinCall(const ast::CallExpression* expr,
     bool has_side_effects =
         builtin.sem->HasSideEffects() ||
         std::any_of(args.begin(), args.end(), [](auto* e) { return e->HasSideEffects(); });
-    auto* call = builder_->create<sem::Call>(expr, builtin.sem, std::move(args), current_statement_,
+    auto* call = builder_->create<sem::Call>(expr, builtin.sem, stage, std::move(args),
-                                             value, has_side_effects);
+                                             current_statement_, value, has_side_effects);
 
     if (current_function_) {
         current_function_->AddDirectlyCalledBuiltin(builtin.sem);
@@ -1755,7 +1785,8 @@ sem::Call* Resolver::FunctionCall(const ast::CallExpression* expr,
     // TODO(crbug.com/tint/1420): For now, assume all function calls have side
     // effects.
     bool has_side_effects = true;
-    auto* call = builder_->create<sem::Call>(expr, target, std::move(args), current_statement_,
+    auto* call = builder_->create<sem::Call>(expr, target, sem::EvaluationStage::kRuntime,
+                                             std::move(args), current_statement_,
                                              /* constant_value */ nullptr, has_side_effects);
 
     target->AddCallSite(call);
@@ -1851,7 +1882,8 @@ sem::Expression* Resolver::Literal(const ast::LiteralExpression* literal) {
     }
 
     auto val = const_eval_.Literal(ty, literal);
-    return builder_->create<sem::Expression>(literal, ty, current_statement_, std::move(val),
+    return builder_->create<sem::Expression>(literal, ty, sem::EvaluationStage::kConstant,
+                                             current_statement_, std::move(val),
                                              /* has_side_effects */ false);
 }
 
@@ -2054,6 +2086,7 @@ sem::Expression* Resolver::Binary(const ast::BinaryExpression* expr) {
     const auto* rhs = sem_.Get(expr->rhs);
     auto* lhs_ty = lhs->Type()->UnwrapRef();
     auto* rhs_ty = rhs->Type()->UnwrapRef();
+    auto stage = sem::EvaluationStage::kRuntime;  // TODO(crbug.com/tint/1581)
 
     auto op = intrinsic_table_->Lookup(expr->op, lhs_ty, rhs_ty, expr->source, false);
     if (!op.result) {
@@ -2079,7 +2112,7 @@ sem::Expression* Resolver::Binary(const ast::BinaryExpression* expr) {
     }
 
     bool has_side_effects = lhs->HasSideEffects() || rhs->HasSideEffects();
-    auto* sem = builder_->create<sem::Expression>(expr, op.result, current_statement_, value,
+    auto* sem = builder_->create<sem::Expression>(expr, op.result, stage, current_statement_, value,
                                                   has_side_effects);
     sem->Behaviors() = lhs->Behaviors() + rhs->Behaviors();
 
@@ -2096,6 +2129,7 @@ sem::Expression* Resolver::UnaryOp(const ast::UnaryOpExpression* unary) {
     const sem::Type* ty = nullptr;
     const sem::Variable* source_var = nullptr;
     const sem::Constant* value = nullptr;
+    auto stage = sem::EvaluationStage::kRuntime;  // TODO(crbug.com/tint/1581)
 
     switch (unary->op) {
         case ast::UnaryOp::kAddressOf:
@@ -2155,7 +2189,7 @@ sem::Expression* Resolver::UnaryOp(const ast::UnaryOpExpression* unary) {
         }
     }
 
-    auto* sem = builder_->create<sem::Expression>(unary, ty, current_statement_, value,
+    auto* sem = builder_->create<sem::Expression>(unary, ty, stage, current_statement_, value,
                                                   expr->HasSideEffects(), source_var);
     sem->Behaviors() = expr->Behaviors();
     return sem;

src/tint/resolver/resolver.h

@@ -400,12 +400,15 @@ class Resolver {
     bool IsBuiltin(Symbol) const;
 
     // ArrayConstructorSig represents a unique array constructor signature.
-    // It is a tuple of the array type and number of arguments provided.
+    // It is a tuple of the array type, number of arguments provided and earliest evaluation stage.
-    using ArrayConstructorSig = utils::UnorderedKeyWrapper<std::tuple<const sem::Array*, size_t>>;
+    using ArrayConstructorSig =
+        utils::UnorderedKeyWrapper<std::tuple<const sem::Array*, size_t, sem::EvaluationStage>>;
 
     // StructConstructorSig represents a unique structure constructor signature.
-    // It is a tuple of the structure type and number of arguments provided.
+    // It is a tuple of the structure type, number of arguments provided and earliest evaluation
-    using StructConstructorSig = utils::UnorderedKeyWrapper<std::tuple<const sem::Struct*, size_t>>;
+    // stage.
+    using StructConstructorSig =
+        utils::UnorderedKeyWrapper<std::tuple<const sem::Struct*, size_t, sem::EvaluationStage>>;
 
     ProgramBuilder* const builder_;
     diag::List& diagnostics_;

src/tint/resolver/validator.cc

@@ -549,7 +549,6 @@ bool Validator::LocalVariable(const sem::Variable* v) const {
             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_)
@@ -567,6 +566,13 @@ bool Validator::GlobalVariable(
     bool ok = Switch(
         decl,  //
         [&](const ast::Var* var) {
+            if (auto* init = global->Constructor();
+                init && init->Stage() > sem::EvaluationStage::kOverride) {
+                AddError("module-scope 'var' initializer must be a constant or override expression",
+                         init->Declaration()->source);
+                return false;
+            }
+
             if (global->StorageClass() == ast::StorageClass::kNone) {
                 AddError("module-scope 'var' declaration must have a storage class", decl->source);
                 return false;
@@ -619,31 +625,7 @@ bool Validator::GlobalVariable(
 
             return Var(global);
         },
-        [&](const ast::Override*) {
+        [&](const ast::Override*) { return Override(global, constant_ids); },
-            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",
@@ -777,10 +759,39 @@ bool Validator::Let(const sem::Variable* v) const {
     return true;
 }
 
-bool Validator::Override(const sem::Variable* v) const {
+bool Validator::Override(const sem::Variable* v,
+                         std::unordered_map<uint32_t, const sem::Variable*> constant_ids) const {
     auto* decl = v->Declaration();
     auto* storage_ty = v->Type()->UnwrapRef();
 
+    if (auto* init = v->Constructor(); init && init->Stage() > sem::EvaluationStage::kOverride) {
+        AddError("'override' initializer must be an override expression",
+                 init->Declaration()->source);
+        return false;
+    }
+
+    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 != v) {
+                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;
+        }
+    }
+
     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'",

src/tint/resolver/validator.h

@@ -371,8 +371,10 @@ class Validator {
 
     /// Validates a 'override' variable declaration
     /// @param v the variable to validate
+    /// @param constant_ids the set of constant ids in the module
     /// @returns true on success, false otherwise.
-    bool Override(const sem::Variable* v) const;
+    bool Override(const sem::Variable* v,
+                  std::unordered_map<uint32_t, const sem::Variable*> constant_ids) const;
 
     /// Validates a 'const' variable declaration
     /// @param v the variable to validate

src/tint/sem/builtin.cc

@@ -90,9 +90,10 @@ bool IsDP4aBuiltin(BuiltinType i) {
 Builtin::Builtin(BuiltinType type,
                  const sem::Type* return_type,
                  std::vector<Parameter*> parameters,
+                 EvaluationStage eval_stage,
                  PipelineStageSet supported_stages,
                  bool is_deprecated)
-    : Base(return_type, utils::ToConstPtrVec(parameters)),
+    : Base(return_type, utils::ToConstPtrVec(parameters), eval_stage),
       type_(type),
       supported_stages_(supported_stages),
       is_deprecated_(is_deprecated) {

src/tint/sem/builtin.h

@@ -83,6 +83,7 @@ class Builtin final : public Castable<Builtin, CallTarget> {
     /// @param type the builtin type
     /// @param return_type the return type for the builtin call
     /// @param parameters the parameters for the builtin overload
+    /// @param eval_stage the earliest evaluation stage for a call to the builtin
     /// @param supported_stages the pipeline stages that this builtin can be
     /// used in
     /// @param is_deprecated true if the particular overload is considered
@@ -90,6 +91,7 @@ class Builtin final : public Castable<Builtin, CallTarget> {
     Builtin(BuiltinType type,
             const sem::Type* return_type,
             std::vector<Parameter*> parameters,
+            EvaluationStage eval_stage,
             PipelineStageSet supported_stages,
             bool is_deprecated);
 

src/tint/sem/call.cc

@@ -23,13 +23,17 @@ namespace tint::sem {
 
 Call::Call(const ast::CallExpression* declaration,
            const CallTarget* target,
+           EvaluationStage stage,
            std::vector<const sem::Expression*> arguments,
            const Statement* statement,
            const Constant* constant,
            bool has_side_effects)
-    : Base(declaration, target->ReturnType(), statement, constant, has_side_effects),
+    : Base(declaration, target->ReturnType(), stage, statement, constant, has_side_effects),
       target_(target),
-      arguments_(std::move(arguments)) {}
+      arguments_(std::move(arguments)) {
+    // Check that the stage is no earlier than the target supports
+    TINT_ASSERT(Semantic, target->Stage() <= stage);
+}
 
 Call::~Call() = default;
 

src/tint/sem/call.h

@@ -30,12 +30,14 @@ class Call final : public Castable<Call, Expression> {
     /// Constructor
     /// @param declaration the AST node
     /// @param target the call target
+    /// @param stage the earliest evaluation stage for the expression
     /// @param arguments the call arguments
     /// @param statement the statement that owns this expression
     /// @param constant the constant value of this expression
     /// @param has_side_effects whether this expression may have side effects
     Call(const ast::CallExpression* declaration,
          const CallTarget* target,
+         EvaluationStage stage,
          std::vector<const sem::Expression*> arguments,
          const Statement* statement,
          const Constant* constant,

src/tint/sem/call_target.cc

@@ -21,8 +21,10 @@ TINT_INSTANTIATE_TYPEINFO(tint::sem::CallTarget);
 
 namespace tint::sem {
 
-CallTarget::CallTarget(const sem::Type* return_type, const ParameterList& parameters)
+CallTarget::CallTarget(const sem::Type* return_type,
-    : signature_{return_type, parameters} {
+                       const ParameterList& parameters,
+                       EvaluationStage stage)
+    : signature_{return_type, parameters}, stage_(stage) {
     TINT_ASSERT(Semantic, return_type);
 }
 

src/tint/sem/call_target.h

@@ -63,9 +63,12 @@ struct CallTargetSignature {
 class CallTarget : public Castable<CallTarget, Node> {
   public:
     /// Constructor
+    /// @param stage the earliest evaluation stage for a call to this target
     /// @param return_type the return type of the call target
     /// @param parameters the parameters for the call target
-    CallTarget(const sem::Type* return_type, const ParameterList& parameters);
+    CallTarget(const sem::Type* return_type,
+               const ParameterList& parameters,
+               EvaluationStage stage);
 
     /// Copy constructor
     CallTarget(const CallTarget&);
@@ -82,8 +85,12 @@ class CallTarget : public Castable<CallTarget, Node> {
     /// @return the signature of the call target
     const CallTargetSignature& Signature() const { return signature_; }
 
+    /// @return the earliest evaluation stage for a call to this target
+    EvaluationStage Stage() const { return stage_; }
+
   private:
     CallTargetSignature signature_;
+    EvaluationStage stage_;
 };
 
 }  // namespace tint::sem

src/tint/sem/evaluation_stage.h

@@ -0,0 +1,60 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_SEM_EVALUATION_STAGE_H_
+#define SRC_TINT_SEM_EVALUATION_STAGE_H_
+
+#include <algorithm>
+#include <initializer_list>
+
+namespace tint::sem {
+
+/// The earliest point in time that an expression can be evaluated
+enum class EvaluationStage {
+    /// Expression can be evaluated at shader creation time
+    kConstant,
+    /// Expression can be evaluated at pipeline creation time
+    kOverride,
+    /// Expression can be evaluated at runtime
+    kRuntime,
+};
+
+/// @returns true if stage `a` comes earlier than stage `b`
+inline bool operator<(EvaluationStage a, EvaluationStage b) {
+    return static_cast<int>(a) < static_cast<int>(b);
+}
+
+/// @returns true if stage `a` comes later than stage `b`
+inline bool operator>(EvaluationStage a, EvaluationStage b) {
+    return static_cast<int>(a) > static_cast<int>(b);
+}
+
+/// @param stages a list of EvaluationStage.
+/// @returns the earliest stage supported by all the provided stages
+inline EvaluationStage EarliestStage(std::initializer_list<EvaluationStage> stages) {
+    auto earliest = EvaluationStage::kConstant;
+    for (auto stage : stages) {
+        earliest = std::max(stage, earliest);
+    }
+    return static_cast<EvaluationStage>(earliest);
+}
+
+template <typename... ARGS>
+inline EvaluationStage EarliestStage(ARGS... args) {
+    return EarliestStage({args...});
+}
+
+}  // namespace tint::sem
+
+#endif  // SRC_TINT_SEM_EVALUATION_STAGE_H_

src/tint/sem/expression.cc

@@ -24,6 +24,7 @@ namespace tint::sem {
 
 Expression::Expression(const ast::Expression* declaration,
                        const sem::Type* type,
+                       EvaluationStage stage,
                        const Statement* statement,
                        const Constant* constant,
                        bool has_side_effects,
@@ -31,10 +32,12 @@ Expression::Expression(const ast::Expression* declaration,
     : declaration_(declaration),
       source_variable_(source_var),
       type_(type),
+      stage_(stage),
       statement_(statement),
       constant_(std::move(constant)),
       has_side_effects_(has_side_effects) {
     TINT_ASSERT(Semantic, type_);
+    TINT_ASSERT(Semantic, (constant != nullptr) == (stage == EvaluationStage::kConstant));
 }
 
 Expression::~Expression() = default;

src/tint/sem/expression.h

@@ -18,6 +18,7 @@
 #include "src/tint/ast/expression.h"
 #include "src/tint/sem/behavior.h"
 #include "src/tint/sem/constant.h"
+#include "src/tint/sem/evaluation_stage.h"
 #include "src/tint/sem/node.h"
 
 // Forward declarations
@@ -28,18 +29,21 @@ class Variable;
 }  // namespace tint::sem
 
 namespace tint::sem {
+
 /// Expression holds the semantic information for expression nodes.
 class Expression : public Castable<Expression, Node> {
   public:
     /// Constructor
     /// @param declaration the AST node
     /// @param type the resolved type of the expression
+    /// @param stage the earliest evaluation stage for the expression
     /// @param statement the statement that owns this expression
     /// @param constant the constant value of the expression. May be null
     /// @param has_side_effects true if this expression may have side-effects
     /// @param source_var the (optional) source variable for this expression
     Expression(const ast::Expression* declaration,
                const sem::Type* type,
+               EvaluationStage stage,
                const Statement* statement,
                const Constant* constant,
                bool has_side_effects,
@@ -54,6 +58,9 @@ class Expression : public Castable<Expression, Node> {
     /// @return the resolved type of the expression
     const sem::Type* Type() const { return type_; }
 
+    /// @return the earliest evaluation stage for the expression
+    EvaluationStage Stage() const { return stage_; }
+
     /// @return the statement that owns this expression
     const Statement* Stmt() const { return statement_; }
 
@@ -87,6 +94,7 @@ class Expression : public Castable<Expression, Node> {
 
   private:
     const sem::Type* const type_;
+    const EvaluationStage stage_;
     const Statement* const statement_;
     const Constant* const constant_;
     sem::Behaviors behaviors_{sem::Behavior::kNext};

src/tint/sem/expression_test.cc

@@ -43,7 +43,8 @@ using ExpressionTest = TestHelper;
 
 TEST_F(ExpressionTest, UnwrapMaterialize) {
     MockConstant c(create<I32>());
-    auto* a = create<Expression>(/* declaration */ nullptr, create<I32>(), /* statement */ nullptr,
+    auto* a = create<Expression>(/* declaration */ nullptr, create<I32>(),
+                                 sem::EvaluationStage::kRuntime, /* statement */ nullptr,
                                  /* constant_value */ nullptr,
                                  /* has_side_effects */ false, /* source_var */ nullptr);
     auto* b = create<Materialize>(a, /* statement */ nullptr, &c);

src/tint/sem/function.cc

@@ -30,7 +30,7 @@ namespace tint::sem {
 Function::Function(const ast::Function* declaration,
                    Type* return_type,
                    std::vector<Parameter*> parameters)
-    : Base(return_type, utils::ToConstPtrVec(parameters)),
+    : Base(return_type, utils::ToConstPtrVec(parameters), EvaluationStage::kRuntime),
       declaration_(declaration),
       workgroup_size_{WorkgroupDimension{1}, WorkgroupDimension{1}, WorkgroupDimension{1}} {
     for (auto* parameter : parameters) {

src/tint/sem/index_accessor_expression.cc

@@ -24,13 +24,14 @@ namespace tint::sem {
 
 IndexAccessorExpression::IndexAccessorExpression(const ast::IndexAccessorExpression* declaration,
                                                  const sem::Type* type,
+                                                 EvaluationStage stage,
                                                  const Expression* object,
                                                  const Expression* index,
                                                  const Statement* statement,
                                                  const Constant* constant,
                                                  bool has_side_effects,
                                                  const Variable* source_var /* = nullptr */)
-    : Base(declaration, type, statement, constant, has_side_effects, source_var),
+    : Base(declaration, type, stage, statement, constant, has_side_effects, source_var),
       object_(object),
       index_(index) {}
 

src/tint/sem/index_accessor_expression.h

@@ -32,6 +32,7 @@ class IndexAccessorExpression final : public Castable<IndexAccessorExpression, E
     /// Constructor
     /// @param declaration the AST node
     /// @param type the resolved type of the expression
+    /// @param stage the earliest evaluation stage for the expression
     /// @param object the object expression that is being indexed
     /// @param index the index expression
     /// @param statement the statement that owns this expression
@@ -40,6 +41,7 @@ class IndexAccessorExpression final : public Castable<IndexAccessorExpression, E
     /// @param source_var the (optional) source variable for this expression
     IndexAccessorExpression(const ast::IndexAccessorExpression* declaration,
                             const sem::Type* type,
+                            EvaluationStage stage,
                             const Expression* object,
                             const Expression* index,
                             const Statement* statement,

src/tint/sem/materialize.cc

@@ -22,6 +22,7 @@ Materialize::Materialize(const Expression* expr,
                          const Constant* constant)
     : Base(/* declaration */ expr->Declaration(),
            /* type */ constant->Type(),
+           /* stage */ EvaluationStage::kConstant,  // Abstract can only be const-expr
            /* statement */ statement,
            /* constant */ constant,
            /* has_side_effects */ false,

src/tint/sem/member_accessor_expression.cc

@@ -25,12 +25,14 @@ namespace tint::sem {
 
 MemberAccessorExpression::MemberAccessorExpression(const ast::MemberAccessorExpression* declaration,
                                                    const sem::Type* type,
+                                                   EvaluationStage stage,
                                                    const Statement* statement,
                                                    const Constant* constant,
                                                    const Expression* object,
                                                    bool has_side_effects,
                                                    const Variable* source_var /* = nullptr */)
-    : Base(declaration, type, statement, constant, has_side_effects, source_var), object_(object) {}
+    : Base(declaration, type, stage, statement, constant, has_side_effects, source_var),
+      object_(object) {}
 
 MemberAccessorExpression::~MemberAccessorExpression() = default;
 
@@ -42,7 +44,14 @@ StructMemberAccess::StructMemberAccess(const ast::MemberAccessorExpression* decl
                                        const StructMember* member,
                                        bool has_side_effects,
                                        const Variable* source_var /* = nullptr */)
-    : Base(declaration, type, statement, constant, object, has_side_effects, source_var),
+    : Base(declaration,
+           type,
+           object->Stage(),
+           statement,
+           constant,
+           object,
+           has_side_effects,
+           source_var),
       member_(member) {}
 
 StructMemberAccess::~StructMemberAccess() = default;
@@ -55,7 +64,14 @@ Swizzle::Swizzle(const ast::MemberAccessorExpression* declaration,
                  std::vector<uint32_t> indices,
                  bool has_side_effects,
                  const Variable* source_var /* = nullptr */)
-    : Base(declaration, type, statement, constant, object, has_side_effects, source_var),
+    : Base(declaration,
+           type,
+           object->Stage(),
+           statement,
+           constant,
+           object,
+           has_side_effects,
+           source_var),
       indices_(std::move(indices)) {}
 
 Swizzle::~Swizzle() = default;

src/tint/sem/member_accessor_expression.h

@@ -37,6 +37,7 @@ class MemberAccessorExpression : public Castable<MemberAccessorExpression, Expre
     /// Constructor
     /// @param declaration the AST node
     /// @param type the resolved type of the expression
+    /// @param stage the earliest evaluation stage for the expression
     /// @param statement the statement that owns this expression
     /// @param constant the constant value of the expression. May be null.
     /// @param object the object that holds the member being accessed
@@ -44,6 +45,7 @@ class MemberAccessorExpression : public Castable<MemberAccessorExpression, Expre
     /// @param source_var the (optional) source variable for this expression
     MemberAccessorExpression(const ast::MemberAccessorExpression* declaration,
                              const sem::Type* type,
+                             EvaluationStage stage,
                              const Statement* statement,
                              const Constant* constant,
                              const Expression* object,

src/tint/sem/type_constructor.cc

@@ -18,8 +18,10 @@ TINT_INSTANTIATE_TYPEINFO(tint::sem::TypeConstructor);
 
 namespace tint::sem {
 
-TypeConstructor::TypeConstructor(const sem::Type* type, const ParameterList& parameters)
+TypeConstructor::TypeConstructor(const sem::Type* type,
-    : Base(type, parameters) {}
+                                 const ParameterList& parameters,
+                                 EvaluationStage stage)
+    : Base(type, parameters, stage) {}
 
 TypeConstructor::~TypeConstructor() = default;
 

src/tint/sem/type_constructor.h

@@ -25,7 +25,8 @@ class TypeConstructor final : public Castable<TypeConstructor, CallTarget> {
     /// Constructor
     /// @param type the type that's being constructed
     /// @param parameters the type constructor parameters
-    TypeConstructor(const sem::Type* type, const ParameterList& parameters);
+    /// @param stage the earliest evaluation stage for the expression
+    TypeConstructor(const sem::Type* type, const ParameterList& parameters, EvaluationStage stage);
 
     /// Destructor
     ~TypeConstructor() override;

src/tint/sem/type_conversion.cc

@@ -18,8 +18,10 @@ TINT_INSTANTIATE_TYPEINFO(tint::sem::TypeConversion);
 
 namespace tint::sem {
 
-TypeConversion::TypeConversion(const sem::Type* type, const sem::Parameter* parameter)
+TypeConversion::TypeConversion(const sem::Type* type,
-    : Base(type, ParameterList{parameter}) {}
+                               const sem::Parameter* parameter,
+                               EvaluationStage stage)
+    : Base(type, ParameterList{parameter}, stage) {}
 
 TypeConversion::~TypeConversion() = default;
 

src/tint/sem/type_conversion.h

@@ -25,7 +25,8 @@ class TypeConversion final : public Castable<TypeConversion, CallTarget> {
     /// Constructor
     /// @param type the target type of the cast
     /// @param parameter the type cast parameter
-    TypeConversion(const sem::Type* type, const sem::Parameter* parameter);
+    /// @param stage the earliest evaluation stage for the expression
+    TypeConversion(const sem::Type* type, const sem::Parameter* parameter, EvaluationStage stage);
 
     /// Destructor
     ~TypeConversion() override;

src/tint/sem/variable.cc

@@ -28,14 +28,15 @@ TINT_INSTANTIATE_TYPEINFO(tint::sem::Parameter);
 TINT_INSTANTIATE_TYPEINFO(tint::sem::VariableUser);
 
 namespace tint::sem {
-
 Variable::Variable(const ast::Variable* declaration,
                    const sem::Type* type,
+                   EvaluationStage stage,
                    ast::StorageClass storage_class,
                    ast::Access access,
                    const Constant* constant_value)
     : declaration_(declaration),
       type_(type),
+      stage_(stage),
       storage_class_(storage_class),
       access_(access),
       constant_value_(constant_value) {}
@@ -44,21 +45,24 @@ Variable::~Variable() = default;
 
 LocalVariable::LocalVariable(const ast::Variable* declaration,
                              const sem::Type* type,
+                             EvaluationStage stage,
                              ast::StorageClass storage_class,
                              ast::Access access,
                              const sem::Statement* statement,
                              const Constant* constant_value)
-    : Base(declaration, type, storage_class, access, constant_value), statement_(statement) {}
+    : Base(declaration, type, stage, storage_class, access, constant_value),
+      statement_(statement) {}
 
 LocalVariable::~LocalVariable() = default;
 
 GlobalVariable::GlobalVariable(const ast::Variable* declaration,
                                const sem::Type* type,
+                               EvaluationStage stage,
                                ast::StorageClass storage_class,
                                ast::Access access,
                                const Constant* constant_value,
                                sem::BindingPoint binding_point)
-    : Base(declaration, type, storage_class, access, constant_value),
+    : Base(declaration, type, stage, storage_class, access, constant_value),
       binding_point_(binding_point) {}
 
 GlobalVariable::~GlobalVariable() = default;
@@ -69,7 +73,9 @@ Parameter::Parameter(const ast::Parameter* declaration,
                      ast::StorageClass storage_class,
                      ast::Access access,
                      const ParameterUsage usage /* = ParameterUsage::kNone */)
-    : Base(declaration, type, storage_class, access, nullptr), index_(index), usage_(usage) {}
+    : Base(declaration, type, EvaluationStage::kRuntime, storage_class, access, nullptr),
+      index_(index),
+      usage_(usage) {}
 
 Parameter::~Parameter() = default;
 
@@ -78,6 +84,7 @@ VariableUser::VariableUser(const ast::IdentifierExpression* declaration,
                            sem::Variable* variable)
     : Base(declaration,
            variable->Type(),
+           variable->Stage(),
            statement,
            variable->ConstantValue(),
            /* has_side_effects */ false),

src/tint/sem/variable.h

@@ -45,11 +45,13 @@ class Variable : public Castable<Variable, Node> {
     /// Constructor
     /// @param declaration the AST declaration node
     /// @param type the variable type
+    /// @param stage the evaluation stage for an expression of this variable type
     /// @param storage_class the variable storage class
     /// @param access the variable access control type
     /// @param constant_value the constant value for the variable. May be null
     Variable(const ast::Variable* declaration,
              const sem::Type* type,
+             EvaluationStage stage,
              ast::StorageClass storage_class,
              ast::Access access,
              const Constant* constant_value);
@@ -63,6 +65,9 @@ class Variable : public Castable<Variable, Node> {
     /// @returns the canonical type for the variable
     const sem::Type* Type() const { return type_; }
 
+    /// @returns the evaluation stage for an expression of this variable type
+    EvaluationStage Stage() const { return stage_; }
+
     /// @returns the storage class for the variable
     ast::StorageClass StorageClass() const { return storage_class_; }
 
@@ -89,6 +94,7 @@ class Variable : public Castable<Variable, Node> {
   private:
     const ast::Variable* const declaration_;
     const sem::Type* const type_;
+    const EvaluationStage stage_;
     const ast::StorageClass storage_class_;
     const ast::Access access_;
     const Constant* constant_value_;
@@ -102,12 +108,14 @@ class LocalVariable final : public Castable<LocalVariable, Variable> {
     /// Constructor
     /// @param declaration the AST declaration node
     /// @param type the variable type
+    /// @param stage the evaluation stage for an expression of this variable type
     /// @param storage_class the variable storage class
     /// @param access the variable access control type
     /// @param statement the statement that declared this local variable
     /// @param constant_value the constant value for the variable. May be null
     LocalVariable(const ast::Variable* declaration,
                   const sem::Type* type,
+                  EvaluationStage stage,
                   ast::StorageClass storage_class,
                   ast::Access access,
                   const sem::Statement* statement,
@@ -137,12 +145,14 @@ class GlobalVariable final : public Castable<GlobalVariable, Variable> {
     /// Constructor
     /// @param declaration the AST declaration node
     /// @param type the variable type
+    /// @param stage the evaluation stage for an expression of this variable type
     /// @param storage_class the variable storage class
     /// @param access the variable access control type
     /// @param constant_value the constant value for the variable. May be null
     /// @param binding_point the optional resource binding point of the variable
     GlobalVariable(const ast::Variable* declaration,
                    const sem::Type* type,
+                   EvaluationStage stage,
                    ast::StorageClass storage_class,
                    ast::Access access,
                    const Constant* constant_value,

src/tint/writer/append_vector.cc

@@ -59,7 +59,8 @@ const sem::Expression* Zero(ProgramBuilder& b, const sem::Type* ty, const sem::S
             << "unsupported vector element type: " << ty->TypeInfo().name;
         return nullptr;
     }
-    auto* sem = b.create<sem::Expression>(expr, ty, stmt, /* constant_value */ nullptr,
+    auto* sem = b.create<sem::Expression>(expr, ty, sem::EvaluationStage::kRuntime, stmt,
+                                          /* constant_value */ nullptr,
                                           /* has_side_effects */ false);
     b.Sem().Add(expr, sem);
     return sem;
@@ -136,10 +137,12 @@ const sem::Call* AppendVector(ProgramBuilder* b,
         auto* scalar_cast_target = b->create<sem::TypeConversion>(
             packed_el_sem_ty,
             b->create<sem::Parameter>(nullptr, 0u, scalar_sem->Type()->UnwrapRef(),
-                                      ast::StorageClass::kNone, ast::Access::kUndefined));
+                                      ast::StorageClass::kNone, ast::Access::kUndefined),
+            sem::EvaluationStage::kRuntime);
         auto* scalar_cast_sem = b->create<sem::Call>(
-            scalar_cast_ast, scalar_cast_target, std::vector<const sem::Expression*>{scalar_sem},
+            scalar_cast_ast, scalar_cast_target, sem::EvaluationStage::kRuntime,
-            statement, /* constant_value */ nullptr, /* has_side_effects */ false);
+            std::vector<const sem::Expression*>{scalar_sem}, statement,
+            /* constant_value */ nullptr, /* has_side_effects */ false);
         b->Sem().Add(scalar_cast_ast, scalar_cast_sem);
         packed.emplace_back(scalar_cast_sem);
     } else {
@@ -151,14 +154,16 @@ const sem::Call* AppendVector(ProgramBuilder* b,
         utils::Transform(packed, [&](const sem::Expression* expr) { return expr->Declaration(); }));
     auto* constructor_target = b->create<sem::TypeConstructor>(
         packed_sem_ty,
-        utils::Transform(
+        utils::Transform(packed,
-            packed, [&](const tint::sem::Expression* arg, size_t i) -> const sem::Parameter* {
+                         [&](const tint::sem::Expression* arg, size_t i) -> const sem::Parameter* {
-                return b->create<sem::Parameter>(nullptr, static_cast<uint32_t>(i),
+                             return b->create<sem::Parameter>(
-                                                 arg->Type()->UnwrapRef(), ast::StorageClass::kNone,
+                                 nullptr, static_cast<uint32_t>(i), arg->Type()->UnwrapRef(),
-                                                 ast::Access::kUndefined);
+                                 ast::StorageClass::kNone, ast::Access::kUndefined);
-            }));
+                         }),
-    auto* constructor_sem = b->create<sem::Call>(constructor_ast, constructor_target, packed,
+        sem::EvaluationStage::kRuntime);
-                                                 statement, /* constant_value */ nullptr,
+    auto* constructor_sem =
+        b->create<sem::Call>(constructor_ast, constructor_target, sem::EvaluationStage::kRuntime,
+                             packed, statement, /* constant_value */ nullptr,
                              /* has_side_effects */ false);
     b->Sem().Add(constructor_ast, constructor_sem);
     return constructor_sem;

src/tint/writer/glsl/generator_impl.cc

@@ -1338,7 +1338,8 @@ bool GeneratorImpl::EmitBarrierCall(std::ostream& out, const sem::Builtin* built
 const ast::Expression* GeneratorImpl::CreateF32Zero(const sem::Statement* stmt) {
     auto* zero = builder_.Expr(0_f);
     auto* f32 = builder_.create<sem::F32>();
-    auto* sem_zero = builder_.create<sem::Expression>(zero, f32, stmt, /* constant_value */ nullptr,
+    auto* sem_zero = builder_.create<sem::Expression>(zero, f32, sem::EvaluationStage::kRuntime,
+                                                      stmt, /* constant_value */ nullptr,
                                                       /* has_side_effects */ false);
     builder_.Sem().Add(zero, sem_zero);
     return zero;

src/tint/writer/hlsl/generator_impl.cc

@@ -2392,7 +2392,8 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
         auto* zero = builder_.Expr(0_i);
         auto* stmt = builder_.Sem().Get(vector)->Stmt();
         builder_.Sem().Add(
-            zero, builder_.create<sem::Expression>(zero, i32, stmt, /* constant_value */ nullptr,
+            zero, builder_.create<sem::Expression>(zero, i32, sem::EvaluationStage::kRuntime, stmt,
+                                                   /* constant_value */ nullptr,
                                                    /* has_side_effects */ false));
         auto* packed = AppendVector(&builder_, vector, zero);
         return EmitExpression(out, packed->Declaration());