diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index b59fa37..fb51327 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1131,7 +1131,6 @@
       "sem/atomic_test.cc",
       "sem/bool_test.cc",
       "sem/builtin_test.cc",
-      "sem/constant_test.cc",
       "sem/depth_multisampled_texture_test.cc",
       "sem/depth_texture_test.cc",
       "sem/expression_test.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 755d4a3..3e7154c 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -810,7 +810,6 @@
     sem/atomic.cc
     sem/bool_test.cc
     sem/builtin_test.cc
-    sem/constant_test.cc
     sem/depth_multisampled_texture_test.cc
     sem/depth_texture_test.cc
     sem/expression_test.cc
diff --git a/src/tint/program.cc b/src/tint/program.cc
index 6722a09..1afbd46 100644
--- a/src/tint/program.cc
+++ b/src/tint/program.cc
@@ -38,6 +38,7 @@
       types_(std::move(program.types_)),
       ast_nodes_(std::move(program.ast_nodes_)),
       sem_nodes_(std::move(program.sem_nodes_)),
+      constant_nodes_(std::move(program.constant_nodes_)),
       ast_(std::move(program.ast_)),
       sem_(std::move(program.sem_)),
       symbols_(std::move(program.symbols_)),
@@ -62,6 +63,7 @@
     types_ = std::move(builder.Types());
     ast_nodes_ = std::move(builder.ASTNodes());
     sem_nodes_ = std::move(builder.SemNodes());
+    constant_nodes_ = std::move(builder.ConstantNodes());
     ast_ = &builder.AST();  // ast::Module is actually a heap allocation.
     sem_ = std::move(builder.Sem());
     symbols_ = std::move(builder.Symbols());
@@ -86,6 +88,7 @@
     types_ = std::move(program.types_);
     ast_nodes_ = std::move(program.ast_nodes_);
     sem_nodes_ = std::move(program.sem_nodes_);
+    constant_nodes_ = std::move(program.constant_nodes_);
     ast_ = std::move(program.ast_);
     sem_ = std::move(program.sem_);
     symbols_ = std::move(program.symbols_);
diff --git a/src/tint/program.h b/src/tint/program.h
index 3230e7e..5fd31dd 100644
--- a/src/tint/program.h
+++ b/src/tint/program.h
@@ -20,6 +20,7 @@
 
 #include "src/tint/ast/function.h"
 #include "src/tint/program_id.h"
+#include "src/tint/sem/constant.h"
 #include "src/tint/sem/info.h"
 #include "src/tint/sem/type_manager.h"
 #include "src/tint/symbol_table.h"
@@ -43,6 +44,9 @@
     /// SemNodeAllocator is an alias to BlockAllocator<sem::Node>
     using SemNodeAllocator = utils::BlockAllocator<sem::Node>;
 
+    /// ConstantAllocator is an alias to BlockAllocator<sem::Constant>
+    using ConstantAllocator = utils::BlockAllocator<sem::Constant>;
+
     /// Constructor
     Program();
 
@@ -160,6 +164,7 @@
     sem::Manager types_;
     ASTNodeAllocator ast_nodes_;
     SemNodeAllocator sem_nodes_;
+    ConstantAllocator constant_nodes_;
     ast::Module* ast_ = nullptr;
     sem::Info sem_;
     SymbolTable symbols_{id_};
diff --git a/src/tint/program_builder.h b/src/tint/program_builder.h
index 5bccb5f..e9b0a80 100644
--- a/src/tint/program_builder.h
+++ b/src/tint/program_builder.h
@@ -88,6 +88,7 @@
 #include "src/tint/program_id.h"
 #include "src/tint/sem/array.h"
 #include "src/tint/sem/bool.h"
+#include "src/tint/sem/constant.h"
 #include "src/tint/sem/depth_texture.h"
 #include "src/tint/sem/external_texture.h"
 #include "src/tint/sem/f16.h"
@@ -163,6 +164,9 @@
     /// SemNodeAllocator is an alias to BlockAllocator<sem::Node>
     using SemNodeAllocator = utils::BlockAllocator<sem::Node>;
 
+    /// ConstantAllocator is an alias to BlockAllocator<sem::Constant>
+    using ConstantAllocator = utils::BlockAllocator<sem::Constant>;
+
     /// Constructor
     ProgramBuilder();
 
@@ -229,6 +233,12 @@
         return sem_nodes_;
     }
 
+    /// @returns a reference to the program's semantic constant storage
+    ConstantAllocator& ConstantNodes() {
+        AssertNotMoved();
+        return constant_nodes_;
+    }
+
     /// @returns a reference to the program's AST root Module
     ast::Module& AST() {
         AssertNotMoved();
@@ -332,9 +342,8 @@
     }
 
     /// Creates a new sem::Node owned by the ProgramBuilder.
-    /// When the ProgramBuilder is destructed, the sem::Node will also be
-    /// destructed.
-    /// @param args the arguments to pass to the type constructor
+    /// When the ProgramBuilder is destructed, the sem::Node will also be destructed.
+    /// @param args the arguments to pass to the constructor
     /// @returns the node pointer
     template <typename T, typename... ARGS>
     traits::EnableIf<traits::IsTypeOrDerived<T, sem::Node> &&
@@ -345,6 +354,16 @@
         return sem_nodes_.Create<T>(std::forward<ARGS>(args)...);
     }
 
+    /// Creates a new sem::Constant owned by the ProgramBuilder.
+    /// When the ProgramBuilder is destructed, the sem::Node will also be destructed.
+    /// @param args the arguments to pass to the constructor
+    /// @returns the node pointer
+    template <typename T, typename... ARGS>
+    traits::EnableIf<traits::IsTypeOrDerived<T, sem::Constant>, T>* create(ARGS&&... args) {
+        AssertNotMoved();
+        return constant_nodes_.Create<T>(std::forward<ARGS>(args)...);
+    }
+
     /// Creates a new sem::Type owned by the ProgramBuilder.
     /// When the ProgramBuilder is destructed, owned ProgramBuilder and the
     /// returned`Type` will also be destructed.
@@ -2747,6 +2766,7 @@
     sem::Manager types_;
     ASTNodeAllocator ast_nodes_;
     SemNodeAllocator sem_nodes_;
+    ConstantAllocator constant_nodes_;
     ast::Module* ast_;
     sem::Info sem_;
     SymbolTable symbols_{id_};
diff --git a/src/tint/resolver/const_eval.h b/src/tint/resolver/const_eval.h
index 89bb3da..3792e35 100644
--- a/src/tint/resolver/const_eval.h
+++ b/src/tint/resolver/const_eval.h
@@ -30,7 +30,9 @@
 namespace tint::resolver::const_eval {
 
 /// Typedef for a constant evaluation function
-using Function = sem::Constant(ProgramBuilder& builder, const sem::Constant* args, size_t num_args);
+using Function = const sem::Constant*(ProgramBuilder& builder,
+                                      sem::Constant const* const* args,
+                                      size_t num_args);
 
 }  // namespace tint::resolver::const_eval
 
diff --git a/src/tint/resolver/materialize_test.cc b/src/tint/resolver/materialize_test.cc
index 0ace1ca..1b935ee 100644
--- a/src/tint/resolver/materialize_test.cc
+++ b/src/tint/resolver/materialize_test.cc
@@ -73,12 +73,61 @@
     return o << "<unknown>";
 }
 
+template <typename CASE>
+class MaterializeTest : public resolver::ResolverTestWithParam<CASE> {
+  protected:
+    using ProgramBuilder::FriendlyName;
+
+    void CheckTypesAndValues(const sem::Expression* expr,
+                             const tint::sem::Type* expected_sem_ty,
+                             const std::variant<AInt, AFloat>& expected_value) {
+        std::visit([&](auto v) { CheckTypesAndValuesImpl(expr, expected_sem_ty, v); },
+                   expected_value);
+    }
+
+  private:
+    template <typename T>
+    void CheckTypesAndValuesImpl(const sem::Expression* expr,
+                                 const tint::sem::Type* expected_sem_ty,
+                                 T expected_value) {
+        EXPECT_TYPE(expr->Type(), expected_sem_ty);
+
+        auto* value = expr->ConstantValue();
+        ASSERT_NE(value, nullptr);
+        EXPECT_TYPE(expr->Type(), value->Type());
+
+        tint::Switch(
+            expected_sem_ty,  //
+            [&](const sem::Vector* v) {
+                for (uint32_t i = 0; i < v->Width(); i++) {
+                    auto* el = value->Index(i);
+                    ASSERT_NE(el, nullptr);
+                    EXPECT_TYPE(el->Type(), v->type());
+                    EXPECT_EQ(std::get<T>(el->Value()), expected_value);
+                }
+            },
+            [&](const sem::Matrix* m) {
+                for (uint32_t c = 0; c < m->columns(); c++) {
+                    auto* column = value->Index(c);
+                    ASSERT_NE(column, nullptr);
+                    EXPECT_TYPE(column->Type(), m->ColumnType());
+                    for (uint32_t r = 0; r < m->rows(); r++) {
+                        auto* el = column->Index(r);
+                        ASSERT_NE(el, nullptr);
+                        EXPECT_TYPE(el->Type(), m->type());
+                        EXPECT_EQ(std::get<T>(el->Value()), expected_value);
+                    }
+                }
+            },
+            [&](Default) { EXPECT_EQ(std::get<T>(value->Value()), expected_value); });
+    }
+};
+
 ////////////////////////////////////////////////////////////////////////////////////////////////////
 // MaterializeAbstractNumericToConcreteType
 // Tests that an abstract-numeric will materialize to the expected concrete type
 ////////////////////////////////////////////////////////////////////////////////////////////////////
 namespace materialize_abstract_numeric_to_concrete_type {
-
 // How should the materialization occur?
 enum class Method {
     // var a : target_type = abstract_expr;
@@ -247,7 +296,7 @@
 }
 
 using MaterializeAbstractNumericToConcreteType =
-    resolver::ResolverTestWithParam<std::tuple<Expectation, Method, Data>>;
+    MaterializeTest<std::tuple<Expectation, Method, Data>>;
 
 TEST_P(MaterializeAbstractNumericToConcreteType, Test) {
     // Once built-in and ops using f16 is properly supported, we'll need to enable this:
@@ -323,30 +372,12 @@
             break;
     }
 
-    auto check_types_and_values = [&](const sem::Expression* expr) {
-        auto* target_sem_ty = data.target_sem_ty(*this);
-
-        EXPECT_TYPE(expr->Type(), target_sem_ty);
-        EXPECT_TYPE(expr->ConstantValue().Type(), target_sem_ty);
-
-        uint32_t num_elems = 0;
-        const sem::Type* target_sem_el_ty = sem::Type::DeepestElementOf(target_sem_ty, &num_elems);
-        EXPECT_TYPE(expr->ConstantValue().ElementType(), target_sem_el_ty);
-        expr->ConstantValue().WithElements([&](auto&& vec) {
-            using VEC_TY = std::decay_t<decltype(vec)>;
-            using EL_TY = typename VEC_TY::value_type;
-            ASSERT_TRUE(std::holds_alternative<EL_TY>(data.materialized_value));
-            VEC_TY expected(num_elems, std::get<EL_TY>(data.materialized_value));
-            EXPECT_EQ(vec, expected);
-        });
-    };
-
     switch (expectation) {
         case Expectation::kMaterialize: {
             ASSERT_TRUE(r()->Resolve()) << r()->error();
             auto* materialize = Sem().Get<sem::Materialize>(abstract_expr);
             ASSERT_NE(materialize, nullptr);
-            check_types_and_values(materialize);
+            CheckTypesAndValues(materialize, data.target_sem_ty(*this), data.materialized_value);
             break;
         }
         case Expectation::kNoMaterialize: {
@@ -354,7 +385,7 @@
             auto* sem = Sem().Get(abstract_expr);
             ASSERT_NE(sem, nullptr);
             EXPECT_FALSE(sem->Is<sem::Materialize>());
-            check_types_and_values(sem);
+            CheckTypesAndValues(sem, data.target_sem_ty(*this), data.materialized_value);
             break;
         }
         case Expectation::kInvalidConversion: {
@@ -414,8 +445,8 @@
 /// Methods that do not materialize
 constexpr Method kNoMaterializeMethods[] = {
     Method::kPhonyAssign,
-    // TODO(crbug.com/tint/1504): Enable once we have abstract overloads of builtins / binary ops:
-    // Method::kBuiltinArg, Method::kBinaryOp,
+    // TODO(crbug.com/tint/1504): Enable once we have abstract overloads of builtins / binary
+    // ops: Method::kBuiltinArg, Method::kBinaryOp,
 };
 INSTANTIATE_TEST_SUITE_P(
     MaterializeScalar,
@@ -703,7 +734,7 @@
 }
 
 using MaterializeAbstractNumericToDefaultType =
-    resolver::ResolverTestWithParam<std::tuple<Expectation, Method, Data>>;
+    MaterializeTest<std::tuple<Expectation, Method, Data>>;
 
 TEST_P(MaterializeAbstractNumericToDefaultType, Test) {
     const auto& param = GetParam();
@@ -751,32 +782,14 @@
             break;
     }
 
-    auto check_types_and_values = [&](const sem::Expression* expr) {
-        auto* expected_sem_ty = data.expected_sem_ty(*this);
-
-        EXPECT_TYPE(expr->Type(), expected_sem_ty);
-        EXPECT_TYPE(expr->ConstantValue().Type(), expected_sem_ty);
-
-        uint32_t num_elems = 0;
-        const sem::Type* expected_sem_el_ty =
-            sem::Type::DeepestElementOf(expected_sem_ty, &num_elems);
-        EXPECT_TYPE(expr->ConstantValue().ElementType(), expected_sem_el_ty);
-        expr->ConstantValue().WithElements([&](auto&& vec) {
-            using VEC_TY = std::decay_t<decltype(vec)>;
-            using EL_TY = typename VEC_TY::value_type;
-            ASSERT_TRUE(std::holds_alternative<EL_TY>(data.materialized_value));
-            VEC_TY expected(num_elems, std::get<EL_TY>(data.materialized_value));
-            EXPECT_EQ(vec, expected);
-        });
-    };
-
     switch (expectation) {
         case Expectation::kMaterialize: {
             ASSERT_TRUE(r()->Resolve()) << r()->error();
             for (auto* expr : abstract_exprs) {
                 auto* materialize = Sem().Get<sem::Materialize>(expr);
                 ASSERT_NE(materialize, nullptr);
-                check_types_and_values(materialize);
+                CheckTypesAndValues(materialize, data.expected_sem_ty(*this),
+                                    data.materialized_value);
             }
             break;
         }
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index e6cad77..9ef9a30 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -365,13 +365,13 @@
 
     sem::Variable* sem = nullptr;
     if (is_global) {
-        sem = builder_->create<sem::GlobalVariable>(v, ty, ast::StorageClass::kNone,
-                                                    ast::Access::kUndefined, sem::Constant{},
-                                                    sem::BindingPoint{});
+        sem = builder_->create<sem::GlobalVariable>(
+            v, ty, ast::StorageClass::kNone, ast::Access::kUndefined, /* constant_value */ nullptr,
+            sem::BindingPoint{});
     } else {
         sem = builder_->create<sem::LocalVariable>(v, ty, ast::StorageClass::kNone,
                                                    ast::Access::kUndefined, current_statement_,
-                                                   sem::Constant{});
+                                                   /* constant_value */ nullptr);
     }
 
     sem->SetConstructor(rhs);
@@ -419,9 +419,9 @@
         return nullptr;
     }
 
-    auto* sem = builder_->create<sem::GlobalVariable>(v, ty, ast::StorageClass::kNone,
-                                                      ast::Access::kUndefined, sem::Constant{},
-                                                      sem::BindingPoint{});
+    auto* sem = builder_->create<sem::GlobalVariable>(
+        v, ty, ast::StorageClass::kNone, ast::Access::kUndefined, /* constant_value */ nullptr,
+        sem::BindingPoint{});
 
     if (auto* id = ast::GetAttribute<ast::IdAttribute>(v->attributes)) {
         sem->SetConstantId(static_cast<uint16_t>(id->value));
@@ -564,11 +564,11 @@
             binding_point = {bp.group->value, bp.binding->value};
         }
         sem = builder_->create<sem::GlobalVariable>(var, var_ty, storage_class, access,
-                                                    sem::Constant{}, binding_point);
+                                                    /* constant_value */ nullptr, binding_point);
 
     } else {
-        sem = builder_->create<sem::LocalVariable>(var, var_ty, storage_class, access,
-                                                   current_statement_, sem::Constant{});
+        sem = builder_->create<sem::LocalVariable>(
+            var, var_ty, storage_class, access, current_statement_, /* constant_value */ nullptr);
     }
 
     sem->SetConstructor(rhs);
@@ -916,7 +916,7 @@
             return false;
         }
 
-        sem::Constant value;
+        const sem::Constant* value = nullptr;
 
         if (auto* user = args[i]->As<sem::VariableUser>()) {
             // We have an variable of a module-scope constant.
@@ -950,12 +950,12 @@
             continue;
         }
         // validator_.Validate and set the default value for this dimension.
-        if (value.Element<AInt>(0).value < 1) {
+        if (value->As<AInt>() < 1) {
             AddError("workgroup_size argument must be at least 1", values[i]->source);
             return false;
         }
 
-        ws[i].value = value.Element<uint32_t>(0);
+        ws[i].value = value->As<uint32_t>();
     }
 
     current_function_->SetWorkgroupSize(std::move(ws));
@@ -1266,7 +1266,8 @@
             [&](const ast::UnaryOpExpression* unary) -> sem::Expression* { return UnaryOp(unary); },
             [&](const ast::PhonyExpression*) -> sem::Expression* {
                 return builder_->create<sem::Expression>(expr, builder_->create<sem::Void>(),
-                                                         current_statement_, sem::Constant{},
+                                                         current_statement_,
+                                                         /* constant_value */ nullptr,
                                                          /* has_side_effects */ false);
             },
             [&](Default) {
@@ -1309,13 +1310,14 @@
                 << ") returned invalid value";
             return nullptr;
         }
-        auto materialized_val = ConvertValue(std::move(expr_val), target_ty, decl->source);
+        auto materialized_val = ConvertValue(expr_val, target_ty, decl->source);
         if (!materialized_val) {
+            // ConvertValue() has already failed and raised an diagnostic error.
             return nullptr;
         }
-        if (!materialized_val->IsValid()) {
+        if (!materialized_val.Get()) {
             TINT_ICE(Resolver, builder_->Diagnostics())
-                << decl->source << "ConvertValue(" << builder_->FriendlyName(expr_val.Type())
+                << decl->source << "ConvertValue(" << builder_->FriendlyName(expr_val->Type())
                 << " -> " << builder_->FriendlyName(target_ty) << ") returned invalid value";
             return nullptr;
         }
@@ -1678,9 +1680,9 @@
     }
 
     // If the builtin is @const, and all arguments have constant values, evaluate the builtin now.
-    sem::Constant constant;
+    const sem::Constant* constant = nullptr;
     if (builtin.const_eval_fn) {
-        std::vector<sem::Constant> values(args.size());
+        std::vector<const sem::Constant*> values(args.size());
         bool is_const = true;  // all arguments have constant values
         for (size_t i = 0; i < values.size(); i++) {
             if (auto v = args[i]->ConstantValue()) {
@@ -1757,7 +1759,7 @@
     // effects.
     bool has_side_effects = true;
     auto* call = builder_->create<sem::Call>(expr, target, std::move(args), current_statement_,
-                                             sem::Constant{}, has_side_effects);
+                                             /* constant_value */ nullptr, has_side_effects);
 
     target->AddCallSite(call);
 
@@ -2226,21 +2228,21 @@
             return nullptr;
         }
 
-        auto count_val = count_sem->ConstantValue();
+        auto* count_val = count_sem->ConstantValue();
         if (!count_val) {
             AddError("array size must evaluate to a constant integer expression",
                      count_expr->source);
             return nullptr;
         }
 
-        if (auto* ty = count_val.Type(); !ty->is_integer_scalar()) {
+        if (auto* ty = count_val->Type(); !ty->is_integer_scalar()) {
             AddError("array size must evaluate to a constant integer expression, but is type '" +
                          builder_->FriendlyName(ty) + "'",
                      count_expr->source);
             return nullptr;
         }
 
-        count = count_val.Element<AInt>(0).value;
+        count = count_val->As<AInt>();
         if (count < 1) {
             AddError("array size (" + std::to_string(count) + ") must be greater than 0",
                      count_expr->source);
diff --git a/src/tint/resolver/resolver.h b/src/tint/resolver/resolver.h
index aab772e..1a45bc4 100644
--- a/src/tint/resolver/resolver.h
+++ b/src/tint/resolver/resolver.h
@@ -209,22 +209,30 @@
     /// These methods are called from the expression resolving methods, and so child-expression
     /// nodes are guaranteed to have been already resolved and any constant values calculated.
     ////////////////////////////////////////////////////////////////////////////////////////////////
-    sem::Constant EvaluateConstantValue(const ast::Expression* expr, const sem::Type* type);
-    sem::Constant EvaluateConstantValue(const ast::IdentifierExpression* ident,
-                                        const sem::Type* type);
-    sem::Constant EvaluateConstantValue(const ast::LiteralExpression* literal,
-                                        const sem::Type* type);
-    sem::Constant EvaluateConstantValue(const ast::CallExpression* call, const sem::Type* type);
-    sem::Constant EvaluateConstantValue(const ast::IndexAccessorExpression* call,
-                                        const sem::Type* type);
+    const sem::Constant* EvaluateConstantValue(const ast::Expression* expr, const sem::Type* type);
+    const sem::Constant* EvaluateConstantValue(const ast::IdentifierExpression* ident,
+                                               const sem::Type* type);
+    const sem::Constant* EvaluateConstantValue(const ast::LiteralExpression* literal,
+                                               const sem::Type* type);
+    const sem::Constant* EvaluateConstantValue(const ast::CallExpression* call,
+                                               const sem::Type* type);
+    const sem::Constant* EvaluateConstantValue(const ast::IndexAccessorExpression* call,
+                                               const sem::Type* type);
 
-    /// The result type of a ConstantEvaluation method. Holds the constant value and a boolean,
-    /// which is true on success, false on an error.
-    using ConstantResult = utils::Result<sem::Constant>;
+    /// The result type of a ConstantEvaluation method.
+    /// Can be one of three distinct values:
+    /// * A non-null sem::Constant pointer. Returned when a expression resolves to a creation time
+    ///   value.
+    /// * A null sem::Constant pointer. Returned when a expression cannot resolve to a creation time
+    ///   value, but is otherwise legal.
+    /// * `utils::Failure`. Returned when there was a resolver error. In this situation the method
+    ///   will have already reported a diagnostic error message, and the caller should abort
+    ///   resolving.
+    using ConstantResult = utils::Result<const sem::Constant*>;
 
     /// Convert the `value` to `target_type`
     /// @return the converted value
-    ConstantResult ConvertValue(const sem::Constant& value,
+    ConstantResult ConvertValue(const sem::Constant* value,
                                 const sem::Type* target_type,
                                 const Source& source);
 
diff --git a/src/tint/resolver/resolver_constants.cc b/src/tint/resolver/resolver_constants.cc
index 74e4327..f6ffb28 100644
--- a/src/tint/resolver/resolver_constants.cc
+++ b/src/tint/resolver/resolver_constants.cc
@@ -14,7 +14,6 @@
 
 #include "src/tint/resolver/resolver.h"
 
-#include <cmath>
 #include <optional>
 
 #include "src/tint/sem/abstract_float.h"
@@ -22,8 +21,6 @@
 #include "src/tint/sem/constant.h"
 #include "src/tint/sem/type_constructor.h"
 #include "src/tint/utils/compiler_macros.h"
-#include "src/tint/utils/map.h"
-#include "src/tint/utils/transform.h"
 
 using namespace tint::number_suffixes;  // NOLINT
 
@@ -31,127 +28,334 @@
 
 namespace {
 
-/// Converts and returns all the element values of `in` to the type `T`, using the converter
-/// function `CONVERTER`.
-/// @param elements_in the vector of elements to be converted
-/// @param converter a function-like with the signature `void(TO&, FROM)`
-/// @returns the elements converted to type T.
-template <typename T, typename ELEMENTS_IN, typename CONVERTER>
-sem::Constant::Elements Transform(const ELEMENTS_IN& elements_in, CONVERTER&& converter) {
-    TINT_BEGIN_DISABLE_WARNING(UNREACHABLE_CODE);
+/// TypeDispatch is a helper for calling the function `f`, passing a single zero-value argument of
+/// the C++ type that corresponds to the sem::Type `type`. For example, calling `TypeDispatch()`
+/// with a type of `sem::I32*` will call the function f with a single argument of `i32(0)`.
+/// @returns the value returned by calling `f`.
+/// @note `type` must be a scalar or abstract numeric type. Other types will not call `f`, and will
+/// return the zero-initialized value of the return type for `f`.
+template <typename F>
+auto TypeDispatch(const sem::Type* type, F&& f) {
+    return Switch(
+        type,                                                     //
+        [&](const sem::AbstractInt*) { return f(AInt(0)); },      //
+        [&](const sem::AbstractFloat*) { return f(AFloat(0)); },  //
+        [&](const sem::I32*) { return f(i32(0)); },               //
+        [&](const sem::U32*) { return f(u32(0)); },               //
+        [&](const sem::F32*) { return f(f32(0)); },               //
+        [&](const sem::F16*) { return f(f16(0)); },               //
+        [&](const sem::Bool*) { return f(static_cast<bool>(0)); });
+}
 
-    return utils::Transform(elements_in, [&](auto value_in) {
-        if constexpr (std::is_same_v<UnwrapNumber<T>, bool>) {
-            return AInt(value_in != 0);
+/// @returns `value` if `T` is not a Number, otherwise ValueOf returns the inner value of the
+/// Number.
+template <typename T>
+inline auto ValueOf(T value) {
+    if constexpr (std::is_same_v<UnwrapNumber<T>, T>) {
+        return value;
+    } else {
+        return value.value;
+    }
+}
+
+/// @returns true if `value` is a positive zero.
+template <typename T>
+inline bool IsPositiveZero(T value) {
+    using N = UnwrapNumber<T>;
+    return Number<N>(value) == Number<N>(0);  // Considers sign bit
+}
+
+/// Constant inherits from sem::Constant to add an private implementation method for conversion.
+struct Constant : public sem::Constant {
+    /// Convert attempts to convert the constant value to the given type. On error, Convert()
+    /// creates a new diagnostic message and returns a Failure.
+    virtual utils::Result<const Constant*> Convert(ProgramBuilder& builder,
+                                                   const sem::Type* target_ty,
+                                                   const Source& source) const = 0;
+};
+
+// Forward declaration
+const Constant* CreateComposite(ProgramBuilder& builder,
+                                const sem::Type* type,
+                                std::vector<const Constant*> elements);
+
+/// Element holds a single scalar or abstract-numeric value.
+/// Element implements the Constant interface.
+template <typename T>
+struct Element : Constant {
+    Element(const sem::Type* t, T v) : type(t), value(v) {}
+    ~Element() override = default;
+    const sem::Type* Type() const override { return type; }
+    std::variant<std::monostate, AInt, AFloat> Value() const override {
+        if constexpr (IsFloatingPoint<UnwrapNumber<T>>) {
+            return static_cast<AFloat>(value);
         } else {
-            T converted{};
-            converter(converted, value_in);
-            if constexpr (IsFloatingPoint<UnwrapNumber<T>>) {
-                return AFloat(converted);
+            return static_cast<AInt>(value);
+        }
+    }
+    const Constant* Index(size_t) const override { return nullptr; }
+    bool AllZero() const override { return IsPositiveZero(value); }
+    bool AnyZero() const override { return IsPositiveZero(value); }
+    bool AllEqual() const override { return true; }
+    size_t Hash() const override { return utils::Hash(type, ValueOf(value)); }
+
+    utils::Result<const Constant*> Convert(ProgramBuilder& builder,
+                                           const sem::Type* target_ty,
+                                           const Source& source) const override {
+        TINT_BEGIN_DISABLE_WARNING(UNREACHABLE_CODE);
+        if (target_ty == type) {
+            // If the types are identical, then no conversion is needed.
+            return this;
+        }
+        bool failed = false;
+        auto* res = TypeDispatch(target_ty, [&](auto zero_to) -> const Constant* {
+            // `T` is the source type, `value` is the source value.
+            // `TO` is the target type.
+            using TO = std::decay_t<decltype(zero_to)>;
+            if constexpr (std::is_same_v<TO, bool>) {
+                // [x -> bool]
+                return builder.create<Element<TO>>(target_ty, !IsPositiveZero(value));
+            } else if constexpr (std::is_same_v<T, bool>) {
+                // [bool -> x]
+                return builder.create<Element<TO>>(target_ty, TO(value ? 1 : 0));
+            } else if (auto conv = CheckedConvert<TO>(value)) {
+                // Conversion success
+                return builder.create<Element<TO>>(target_ty, conv.Get());
+                // --- Below this point are the failure cases ---
+            } else if constexpr (std::is_same_v<T, AInt> || std::is_same_v<T, AFloat>) {
+                // [abstract-numeric -> x] - materialization failure
+                std::stringstream ss;
+                ss << "value " << value << " cannot be represented as ";
+                ss << "'" << builder.FriendlyName(target_ty) << "'";
+                builder.Diagnostics().add_error(tint::diag::System::Resolver, ss.str(), source);
+                failed = true;
+            } else if constexpr (IsFloatingPoint<UnwrapNumber<TO>>) {
+                // [x -> floating-point] - number not exactly representable
+                // https://www.w3.org/TR/WGSL/#floating-point-conversion
+                constexpr auto kInf = std::numeric_limits<double>::infinity();
+                switch (conv.Failure()) {
+                    case ConversionFailure::kExceedsNegativeLimit:
+                        return builder.create<Element<TO>>(target_ty, TO(-kInf));
+                    case ConversionFailure::kExceedsPositiveLimit:
+                        return builder.create<Element<TO>>(target_ty, TO(kInf));
+                }
             } else {
-                return AInt(converted);
+                // [x -> integer] - number not exactly representable
+                // https://www.w3.org/TR/WGSL/#floating-point-conversion
+                switch (conv.Failure()) {
+                    case ConversionFailure::kExceedsNegativeLimit:
+                        return builder.create<Element<TO>>(target_ty, TO(TO::kLowest));
+                    case ConversionFailure::kExceedsPositiveLimit:
+                        return builder.create<Element<TO>>(target_ty, TO(TO::kHighest));
+                }
             }
+            return nullptr;  // Expression is not constant.
+        });
+        if (failed) {
+            // A diagnostic error has been raised, and resolving should abort.
+            return utils::Failure;
         }
-    });
-
-    TINT_END_DISABLE_WARNING(UNREACHABLE_CODE);
-}
-
-/// Converts and returns all the element values of `in` to the semantic type `el_ty`, using the
-/// converter function `CONVERTER`.
-/// @param in the constant to convert
-/// @param el_ty the target element type
-/// @param converter a function-like with the signature `void(TO&, FROM)`
-/// @returns the elements converted to `el_ty`
-template <typename CONVERTER>
-sem::Constant::Elements Transform(const sem::Constant::Elements& in,
-                                  const sem::Type* el_ty,
-                                  CONVERTER&& converter) {
-    return std::visit(
-        [&](auto&& v) {
-            return Switch(
-                el_ty,  //
-                [&](const sem::AbstractInt*) { return Transform<AInt>(v, converter); },
-                [&](const sem::AbstractFloat*) { return Transform<AFloat>(v, converter); },
-                [&](const sem::I32*) { return Transform<i32>(v, converter); },
-                [&](const sem::U32*) { return Transform<u32>(v, converter); },
-                [&](const sem::F32*) { return Transform<f32>(v, converter); },
-                [&](const sem::F16*) { return Transform<f16>(v, converter); },
-                [&](const sem::Bool*) { return Transform<bool>(v, converter); },
-                [&](Default) -> sem::Constant::Elements {
-                    diag::List diags;
-                    TINT_UNREACHABLE(Semantic, diags)
-                        << "invalid element type " << el_ty->TypeInfo().name;
-                    return {};
-                });
-        },
-        in);
-}
-
-/// Converts and returns all the elements in `in` to the type `el_ty`.
-/// If the value does not fit in the target type, and:
-///  * the target type is an integer type, then the resulting value will be clamped to the integer's
-///    highest or lowest value.
-///  * the target type is an float type, then the resulting value will be either positive or
-///    negative infinity, based on the sign of the input value.
-/// @param in the input elements
-/// @param el_ty the target element type
-/// @returns the elements converted to `el_ty`
-sem::Constant::Elements ConvertElements(const sem::Constant::Elements& in, const sem::Type* el_ty) {
-    return Transform(in, el_ty, [](auto& el_out, auto el_in) {
-        using OUT = std::decay_t<decltype(el_out)>;
-        if (auto conv = CheckedConvert<OUT>(el_in)) {
-            el_out = conv.Get();
-        } else {
-            constexpr auto kInf = std::numeric_limits<double>::infinity();
-            switch (conv.Failure()) {
-                case ConversionFailure::kExceedsNegativeLimit:
-                    el_out = IsFloatingPoint<UnwrapNumber<OUT>> ? OUT(-kInf) : OUT::kLowest;
-                    break;
-                case ConversionFailure::kExceedsPositiveLimit:
-                    el_out = IsFloatingPoint<UnwrapNumber<OUT>> ? OUT(kInf) : OUT::kHighest;
-                    break;
-            }
-        }
-    });
-}
-
-/// Converts and returns all the elements in `in` to the type `el_ty`, by performing a
-/// `CheckedConvert` on each element value. A single error diagnostic will be raised if an element
-/// value cannot be represented by the target type.
-/// @param in the input elements
-/// @param el_ty the target element type
-/// @returns the elements converted to `el_ty`, or a Failure if some elements could not be
-/// represented by the target type.
-utils::Result<sem::Constant::Elements> MaterializeElements(const sem::Constant::Elements& in,
-                                                           const sem::Type* el_ty,
-                                                           ProgramBuilder& builder,
-                                                           Source source) {
-    std::optional<std::string> failure;
-
-    auto out = Transform(in, el_ty, [&](auto& el_out, auto el_in) {
-        using OUT = std::decay_t<decltype(el_out)>;
-        if (auto conv = CheckedConvert<OUT>(el_in)) {
-            el_out = conv.Get();
-        } else if (!failure.has_value()) {
-            std::stringstream ss;
-            ss << "value " << el_in << " cannot be represented as ";
-            ss << "'" << builder.FriendlyName(el_ty) << "'";
-            failure = ss.str();
-        }
-    });
-
-    if (failure.has_value()) {
-        builder.Diagnostics().add_error(diag::System::Resolver, std::move(failure.value()), source);
-        return utils::Failure;
+        return res;
+        TINT_END_DISABLE_WARNING(UNREACHABLE_CODE);
     }
 
-    return out;
+    sem::Type const* const type;
+    const T value;
+};
+
+/// Splat holds a single Constant value, duplicated as all children.
+/// Splat is used for zero-initializers, 'splat' constructors, or constructors where each element is
+/// identical. Splat may be of a vector, matrix or array type.
+/// Splat implements the Constant interface.
+struct Splat : Constant {
+    Splat(const sem::Type* t, const Constant* e, size_t n) : type(t), el(e), count(n) {}
+    ~Splat() override = default;
+    const sem::Type* Type() const override { return type; }
+    std::variant<std::monostate, AInt, AFloat> Value() const override { return {}; }
+    const Constant* Index(size_t i) const override { return i < count ? el : nullptr; }
+    bool AllZero() const override { return el->AllZero(); }
+    bool AnyZero() const override { return el->AnyZero(); }
+    bool AllEqual() const override { return true; }
+    size_t Hash() const override { return utils::Hash(type, el->Hash(), count); }
+
+    utils::Result<const Constant*> Convert(ProgramBuilder& builder,
+                                           const sem::Type* target_ty,
+                                           const Source& source) const override {
+        // Convert the single splatted element type.
+        auto conv_el = el->Convert(builder, sem::Type::ElementOf(target_ty), source);
+        if (!conv_el) {
+            return utils::Failure;
+        }
+        if (!conv_el.Get()) {
+            return nullptr;
+        }
+        return builder.create<Splat>(target_ty, conv_el.Get(), count);
+    }
+
+    sem::Type const* const type;
+    const Constant* el;
+    const size_t count;
+};
+
+/// Composite holds a number of mixed child Constant values.
+/// Composite may be of a vector, matrix or array type.
+/// If each element is the same type and value, then a Splat would be a more efficient constant
+/// implementation. Use CreateComposite() to create the appropriate Constant type.
+/// Composite implements the Constant interface.
+struct Composite : Constant {
+    Composite(const sem::Type* t, std::vector<const Constant*> els, bool all_0, bool any_0)
+        : type(t), elements(std::move(els)), all_zero(all_0), any_zero(any_0), hash(CalcHash()) {}
+    ~Composite() override = default;
+    const sem::Type* Type() const override { return type; }
+    std::variant<std::monostate, AInt, AFloat> Value() const override { return {}; }
+    const Constant* Index(size_t i) const override {
+        return i < elements.size() ? elements[i] : nullptr;
+    }
+    bool AllZero() const override { return all_zero; }
+    bool AnyZero() const override { return any_zero; }
+    bool AllEqual() const override { return false; /* otherwise this should be a Splat */ }
+    size_t Hash() const override { return hash; }
+
+    utils::Result<const Constant*> Convert(ProgramBuilder& builder,
+                                           const sem::Type* target_ty,
+                                           const Source& source) const override {
+        // Convert each of the composite element types.
+        auto* el_ty = sem::Type::ElementOf(target_ty);
+        std::vector<const Constant*> conv_els;
+        conv_els.reserve(elements.size());
+        for (auto* el : elements) {
+            auto conv_el = el->Convert(builder, el_ty, source);
+            if (!conv_el) {
+                return utils::Failure;
+            }
+            if (!conv_el.Get()) {
+                return nullptr;
+            }
+            conv_els.emplace_back(conv_el.Get());
+        }
+        return CreateComposite(builder, target_ty, std::move(conv_els));
+    }
+
+    size_t CalcHash() {
+        auto h = utils::Hash(type, all_zero, any_zero);
+        for (auto* el : elements) {
+            utils::HashCombine(&h, el->Hash());
+        }
+        return h;
+    }
+
+    sem::Type const* const type;
+    const std::vector<const Constant*> elements;
+    const bool all_zero;
+    const bool any_zero;
+    const size_t hash;
+};
+
+/// CreateElement constructs and returns an Element<T>.
+template <typename T>
+const Constant* CreateElement(ProgramBuilder& builder, const sem::Type* t, T v) {
+    return builder.create<Element<T>>(t, v);
+}
+
+/// ZeroValue returns a Constant for the zero-value of the type `type`.
+const Constant* ZeroValue(ProgramBuilder& builder, const sem::Type* type) {
+    return Switch(
+        type,  //
+        [&](const sem::Vector* v) -> const Constant* {
+            auto* zero_el = ZeroValue(builder, v->type());
+            return builder.create<Splat>(type, zero_el, v->Width());
+        },
+        [&](const sem::Matrix* m) -> const Constant* {
+            auto* zero_el = ZeroValue(builder, m->ColumnType());
+            return builder.create<Splat>(type, zero_el, m->columns());
+        },
+        [&](const sem::Array* a) -> const Constant* {
+            if (auto* zero_el = ZeroValue(builder, a->ElemType())) {
+                return builder.create<Splat>(type, zero_el, a->Count());
+            }
+            return nullptr;
+        },
+        [&](Default) -> const Constant* {
+            return TypeDispatch(type, [&](auto zero) -> const Constant* {
+                return CreateElement(builder, type, zero);
+            });
+        });
+}
+
+/// Equal returns true if the constants `a` and `b` are of the same type and value.
+bool Equal(const sem::Constant* a, const sem::Constant* b) {
+    if (a->Hash() != b->Hash()) {
+        return false;
+    }
+    if (a->Type() != b->Type()) {
+        return false;
+    }
+    return Switch(
+        a->Type(),  //
+        [&](const sem::Vector* vec) {
+            for (size_t i = 0; i < vec->Width(); i++) {
+                if (!Equal(a->Index(i), b->Index(i))) {
+                    return false;
+                }
+            }
+            return true;
+        },
+        [&](const sem::Matrix* mat) {
+            for (size_t i = 0; i < mat->columns(); i++) {
+                if (!Equal(a->Index(i), b->Index(i))) {
+                    return false;
+                }
+            }
+            return true;
+        },
+        [&](const sem::Array* arr) {
+            for (size_t i = 0; i < arr->Count(); i++) {
+                if (!Equal(a->Index(i), b->Index(i))) {
+                    return false;
+                }
+            }
+            return true;
+        },
+        [&](Default) { return a->Value() == b->Value(); });
+}
+
+/// CreateComposite is used to construct a constant of a vector, matrix or array type.
+/// CreateComposite examines the element values and will return either a Composite or a Splat,
+/// depending on the element types and values.
+const Constant* CreateComposite(ProgramBuilder& builder,
+                                const sem::Type* type,
+                                std::vector<const Constant*> elements) {
+    if (elements.size() == 0) {
+        return nullptr;
+    }
+    bool any_zero = false;
+    bool all_zero = true;
+    bool all_equal = true;
+    auto* first = elements.front();
+    for (auto* el : elements) {
+        if (!any_zero && el->AnyZero()) {
+            any_zero = true;
+        }
+        if (all_zero && !el->AllZero()) {
+            all_zero = false;
+        }
+        if (all_equal && el != first) {
+            if (!Equal(el, first)) {
+                all_equal = false;
+            }
+        }
+    }
+    if (all_equal) {
+        return builder.create<Splat>(type, elements[0], elements.size());
+    } else {
+        return builder.create<Composite>(type, std::move(elements), all_zero, any_zero);
+    }
 }
 
 }  // namespace
 
-sem::Constant Resolver::EvaluateConstantValue(const ast::Expression* expr, const sem::Type* type) {
+const sem::Constant* Resolver::EvaluateConstantValue(const ast::Expression* expr,
+                                                     const sem::Type* type) {
     return Switch(
         expr,  //
         [&](const ast::IdentifierExpression* e) { return EvaluateConstantValue(e, type); },
@@ -160,112 +364,176 @@
         [&](const ast::IndexAccessorExpression* e) { return EvaluateConstantValue(e, type); });
 }
 
-sem::Constant Resolver::EvaluateConstantValue(const ast::IdentifierExpression* ident,
-                                              const sem::Type*) {
+const sem::Constant* Resolver::EvaluateConstantValue(const ast::IdentifierExpression* ident,
+                                                     const sem::Type*) {
     if (auto* sem = builder_->Sem().Get(ident)) {
         return sem->ConstantValue();
     }
     return {};
 }
 
-sem::Constant Resolver::EvaluateConstantValue(const ast::LiteralExpression* literal,
-                                              const sem::Type* type) {
+const sem::Constant* Resolver::EvaluateConstantValue(const ast::LiteralExpression* literal,
+                                                     const sem::Type* type) {
     return Switch(
         literal,
         [&](const ast::BoolLiteralExpression* lit) {
-            return sem::Constant{type, {AInt(lit->value ? 1 : 0)}};
+            return CreateElement(*builder_, type, lit->value);
         },
-        [&](const ast::IntLiteralExpression* lit) {
-            return sem::Constant{type, {AInt(lit->value)}};
+        [&](const ast::IntLiteralExpression* lit) -> const Constant* {
+            switch (lit->suffix) {
+                case ast::IntLiteralExpression::Suffix::kNone:
+                    return CreateElement(*builder_, type, AInt(lit->value));
+                case ast::IntLiteralExpression::Suffix::kI:
+                    return CreateElement(*builder_, type, i32(lit->value));
+                case ast::IntLiteralExpression::Suffix::kU:
+                    return CreateElement(*builder_, type, u32(lit->value));
+            }
+            return nullptr;
         },
-        [&](const ast::FloatLiteralExpression* lit) {
-            return sem::Constant{type, {AFloat(lit->value)}};
+        [&](const ast::FloatLiteralExpression* lit) -> const Constant* {
+            switch (lit->suffix) {
+                case ast::FloatLiteralExpression::Suffix::kNone:
+                    return CreateElement(*builder_, type, AFloat(lit->value));
+                case ast::FloatLiteralExpression::Suffix::kF:
+                    return CreateElement(*builder_, type, f32(lit->value));
+                case ast::FloatLiteralExpression::Suffix::kH:
+                    return CreateElement(*builder_, type, f16(lit->value));
+            }
+            return nullptr;
         });
 }
 
-sem::Constant Resolver::EvaluateConstantValue(const ast::CallExpression* call,
-                                              const sem::Type* ty) {
-    uint32_t num_elems = 0;
-    auto* el_ty = sem::Type::DeepestElementOf(ty, &num_elems);
-    if (!el_ty || num_elems == 0) {
-        return {};
-    }
-
+const sem::Constant* Resolver::EvaluateConstantValue(const ast::CallExpression* call,
+                                                     const sem::Type* ty) {
     // Note: we are building constant values for array types. The working group as verbally agreed
     // to support constant expression arrays, but this is not (yet) part of the spec.
     // See: https://github.com/gpuweb/gpuweb/issues/3056
 
     // For zero value init, return 0s
     if (call->args.empty()) {
-        return Switch(
-            el_ty,
-            [&](const sem::AbstractInt*) {
-                return sem::Constant(ty, std::vector(num_elems, AInt(0)));
-            },
-            [&](const sem::AbstractFloat*) {
-                return sem::Constant(ty, std::vector(num_elems, AFloat(0)));
-            },
-            [&](const sem::I32*) { return sem::Constant(ty, std::vector(num_elems, AInt(0))); },
-            [&](const sem::U32*) { return sem::Constant(ty, std::vector(num_elems, AInt(0))); },
-            [&](const sem::F32*) { return sem::Constant(ty, std::vector(num_elems, AFloat(0))); },
-            [&](const sem::F16*) { return sem::Constant(ty, std::vector(num_elems, AFloat(0))); },
-            [&](const sem::Bool*) { return sem::Constant(ty, std::vector(num_elems, AInt(0))); });
+        return ZeroValue(*builder_, ty);
     }
 
-    // Build value for type_ctor from each child value by converting to type_ctor's type.
-    std::optional<sem::Constant::Elements> elements;
-    for (auto* expr : call->args) {
-        auto* arg = builder_->Sem().Get(expr);
+    uint32_t el_count = 0;
+    auto* el_ty = sem::Type::ElementOf(ty, &el_count);
+    if (!el_ty) {
+        return nullptr;  // Target type does not support constant values
+    }
+
+    // value_of returns a `const Constant*` for the expression `expr`, or nullptr if the expression
+    // does not have a constant value.
+    auto value_of = [&](const ast::Expression* expr) {
+        return static_cast<const Constant*>(builder_->Sem().Get(expr)->ConstantValue());
+    };
+
+    if (call->args.size() == 1) {
+        // Type constructor or conversion that takes a single argument.
+        auto& src = call->args[0]->source;
+        auto* arg = value_of(call->args[0]);
         if (!arg) {
-            return {};
-        }
-        auto value = arg->ConstantValue();
-        if (!value) {
-            return {};
+            return nullptr;  // Single argument is not constant.
         }
 
-        // Convert the elements to the desired type.
-        auto converted = ConvertElements(value.GetElements(), el_ty);
-
-        if (elements.has_value()) {
-            // Append the converted vector to elements
-            std::visit(
-                [&](auto&& dst) {
-                    using VEC_TY = std::decay_t<decltype(dst)>;
-                    const auto& src = std::get<VEC_TY>(converted);
-                    dst.insert(dst.end(), src.begin(), src.end());
-                },
-                elements.value());
-        } else {
-            elements = std::move(converted);
+        if (ty->is_scalar()) {  // Scalar type conversion: i32(x), u32(x), bool(x), etc
+            return ConvertValue(arg, el_ty, src).Get();
         }
+
+        if (arg->Type() == el_ty) {
+            // Argument type matches function type. This is a splat.
+            auto splat = [&](size_t n) { return builder_->create<Splat>(ty, arg, n); };
+            return Switch(
+                ty,  //
+                [&](const sem::Vector* v) { return splat(v->Width()); },
+                [&](const sem::Matrix* m) { return splat(m->columns()); },
+                [&](const sem::Array* a) { return splat(a->Count()); });
+        }
+
+        // Argument type and function type mismatch. This is a type conversion.
+        if (auto conv = ConvertValue(arg, ty, src)) {
+            return conv.Get();
+        }
+
+        return nullptr;
     }
 
-    if (!elements) {
-        return {};
-    }
+    // Multiple arguments. Must be a type constructor.
 
-    return std::visit(
-        [&](auto&& v) {
-            if (num_elems != v.size()) {
-                if (v.size() == 1) {
-                    // Splat single-value initializers
-                    for (uint32_t i = 0; i < num_elems - 1; ++i) {
-                        v.emplace_back(v[0]);
+    std::vector<const Constant*> els;  // The constant elements for the composite constant.
+    els.reserve(el_count);
+
+    // Helper for pushing all the argument constants to `els`.
+    auto push_all_args = [&] {
+        for (auto* expr : call->args) {
+            auto* arg = value_of(expr);
+            if (!arg) {
+                return;
+            }
+            els.emplace_back(arg);
+        }
+    };
+
+    Switch(
+        ty,  // What's the target type being constructed?
+        [&](const sem::Vector*) {
+            // Vector can be constructed with a mix of scalars / abstract numerics and smaller
+            // vectors.
+            for (auto* expr : call->args) {
+                auto* arg = value_of(expr);
+                if (!arg) {
+                    return;
+                }
+                auto* arg_ty = arg->Type();
+                if (auto* arg_vec = arg_ty->As<sem::Vector>()) {
+                    // Extract out vector elements.
+                    for (uint32_t i = 0; i < arg_vec->Width(); i++) {
+                        auto* el = static_cast<const Constant*>(arg->Index(i));
+                        if (!el) {
+                            return;
+                        }
+                        els.emplace_back(el);
                     }
                 } else {
-                    // Provided number of arguments does not match the required number of elements.
-                    // Validation should error here.
-                    return sem::Constant{};
+                    els.emplace_back(arg);
                 }
             }
-            return sem::Constant(ty, std::move(elements.value()));
         },
-        elements.value());
+        [&](const sem::Matrix* m) {
+            // Matrix can be constructed with a set of scalars / abstract numerics, or column
+            // vectors.
+            if (call->args.size() == m->columns() * m->rows()) {
+                // Matrix built from scalars / abstract numerics
+                for (uint32_t c = 0; c < m->columns(); c++) {
+                    std::vector<const Constant*> column;
+                    column.reserve(m->rows());
+                    for (uint32_t r = 0; r < m->rows(); r++) {
+                        auto* arg = value_of(call->args[r + c * m->rows()]);
+                        if (!arg) {
+                            return;
+                        }
+                        column.emplace_back(arg);
+                    }
+                    els.push_back(CreateComposite(*builder_, m->ColumnType(), std::move(column)));
+                }
+            } else if (call->args.size() == m->columns()) {
+                // Matrix built from column vectors
+                push_all_args();
+            }
+        },
+        [&](const sem::Array*) {
+            // Arrays must be constructed using a list of elements
+            push_all_args();
+        });
+
+    if (els.size() != el_count) {
+        // If the number of constant elements doesn't match the type, then something went wrong.
+        return nullptr;
+    }
+    // Construct and return either a Composite or Splat.
+    return CreateComposite(*builder_, ty, std::move(els));
 }
 
-sem::Constant Resolver::EvaluateConstantValue(const ast::IndexAccessorExpression* accessor,
-                                              const sem::Type* el_ty) {
+const sem::Constant* Resolver::EvaluateConstantValue(const ast::IndexAccessorExpression* accessor,
+                                                     const sem::Type*) {
     auto* obj_sem = builder_->Sem().Get(accessor->object);
     if (!obj_sem) {
         return {};
@@ -282,20 +550,14 @@
     }
 
     auto idx_val = idx_sem->ConstantValue();
-    if (!idx_val || idx_val.ElementCount() != 1) {
+    if (!idx_val) {
         return {};
     }
 
-    AInt idx = idx_val.Element<AInt>(0);
-
-    // The immediate child element count.
     uint32_t el_count = 0;
-    sem::Type::ElementOf(obj_val.Type(), &el_count);
+    sem::Type::ElementOf(obj_val->Type(), &el_count);
 
-    // The total number of most-nested elements per child element type.
-    uint32_t step = 0;
-    sem::Type::DeepestElementOf(el_ty, &step);
-
+    AInt idx = idx_val->As<AInt>();
     if (idx < 0 || idx >= el_count) {
         auto clamped = std::min<AInt::type>(std::max<AInt::type>(idx, 0), el_count - 1);
         AddWarning("index " + std::to_string(idx) + " out of bounds [0.." +
@@ -305,32 +567,20 @@
         idx = clamped;
     }
 
-    return sem::Constant{el_ty, obj_val.WithElements([&](auto&& v) {
-                             using VEC = std::decay_t<decltype(v)>;
-                             return sem::Constant::Elements(
-                                 VEC(v.begin() + (idx * step), v.begin() + (idx + 1) * step));
-                         })};
+    return obj_val->Index(static_cast<size_t>(idx));
 }
 
-utils::Result<sem::Constant> Resolver::ConvertValue(const sem::Constant& value,
-                                                    const sem::Type* ty,
-                                                    const Source& source) {
-    if (value.Type() == ty) {
+utils::Result<const sem::Constant*> Resolver::ConvertValue(const sem::Constant* value,
+                                                           const sem::Type* target_ty,
+                                                           const Source& source) {
+    if (value->Type() == target_ty) {
         return value;
     }
-
-    auto* el_ty = sem::Type::DeepestElementOf(ty);
-    if (el_ty == nullptr) {
-        return sem::Constant{};
+    auto conv = static_cast<const Constant*>(value)->Convert(*builder_, target_ty, source);
+    if (!conv) {
+        return utils::Failure;
     }
-    if (value.ElementType() == el_ty) {
-        return sem::Constant(ty, value.GetElements());
-    }
-
-    if (auto res = MaterializeElements(value.GetElements(), el_ty, *builder_, source)) {
-        return sem::Constant(ty, std::move(res.Get()));
-    }
-    return utils::Failure;
+    return conv.Get();
 }
 
 }  // namespace tint::resolver
diff --git a/src/tint/resolver/resolver_constants_test.cc b/src/tint/resolver/resolver_constants_test.cc
index 00b886d..ff8a189 100644
--- a/src/tint/resolver/resolver_constants_test.cc
+++ b/src/tint/resolver/resolver_constants_test.cc
@@ -14,10 +14,13 @@
 
 #include "src/tint/resolver/resolver.h"
 
+#include <cmath>
+
 #include "gtest/gtest.h"
 #include "src/tint/resolver/resolver_test_helper.h"
 #include "src/tint/sem/expression.h"
 #include "src/tint/sem/index_accessor_expression.h"
+#include "src/tint/sem/test_helper.h"
 
 using namespace tint::number_suffixes;  // NOLINT
 
@@ -39,10 +42,11 @@
     auto* sem = Sem().Get(expr);
     ASSERT_NE(sem, nullptr);
     EXPECT_TRUE(sem->Type()->Is<sem::I32>());
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_EQ(sem->ConstantValue().ElementType(), sem->Type());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 1u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 99);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->As<AInt>(), 99);
 }
 
 TEST_F(ResolverConstantsTest, Scalar_u32) {
@@ -54,10 +58,11 @@
     auto* sem = Sem().Get(expr);
     ASSERT_NE(sem, nullptr);
     EXPECT_TRUE(sem->Type()->Is<sem::U32>());
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_EQ(sem->ConstantValue().ElementType(), sem->Type());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 1u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 99u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->As<AInt>(), 99u);
 }
 
 TEST_F(ResolverConstantsTest, Scalar_f32) {
@@ -69,10 +74,11 @@
     auto* sem = Sem().Get(expr);
     ASSERT_NE(sem, nullptr);
     EXPECT_TRUE(sem->Type()->Is<sem::F32>());
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_EQ(sem->ConstantValue().ElementType(), sem->Type());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 1u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 9.9f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->As<AFloat>().value, 9.9f);
 }
 
 TEST_F(ResolverConstantsTest, Scalar_f16) {
@@ -85,11 +91,12 @@
     auto* sem = Sem().Get(expr);
     EXPECT_NE(sem, nullptr);
     EXPECT_TRUE(sem->Type()->Is<sem::F16>());
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_EQ(sem->ConstantValue().ElementType(), sem->Type());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 1u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
     // 9.9 is not exactly representable by f16, and should be quantized to 9.8984375
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 9.8984375f);
+    EXPECT_EQ(sem->ConstantValue()->As<AFloat>(), 9.8984375f);
 }
 
 TEST_F(ResolverConstantsTest, Scalar_bool) {
@@ -101,10 +108,11 @@
     auto* sem = Sem().Get(expr);
     ASSERT_NE(sem, nullptr);
     EXPECT_TRUE(sem->Type()->Is<sem::Bool>());
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_EQ(sem->ConstantValue().ElementType(), sem->Type());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 1u);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(0), true);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->As<bool>(), true);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_ZeroInit_i32) {
@@ -119,12 +127,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::I32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 0);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, 0);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, 0);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AInt>(), 0);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AInt>(), 0);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AInt>(), 0);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_ZeroInit_u32) {
@@ -139,12 +160,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::U32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::U32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 0u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, 0u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, 0u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AInt>(), 0u);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AInt>(), 0u);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AInt>(), 0u);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_ZeroInit_f32) {
@@ -159,12 +193,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 0.0);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 0.0);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 0.0);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AFloat>(), 0._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AFloat>(), 0._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AFloat>(), 0._a);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_ZeroInit_f16) {
@@ -180,12 +227,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F16>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 0.0);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 0.0);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 0.0);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AFloat>(), 0._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AFloat>(), 0._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AFloat>(), 0._a);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_ZeroInit_bool) {
@@ -200,12 +260,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::Bool>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::Bool>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(0), false);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(1), false);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(2), false);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<bool>(), false);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<bool>(), false);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<bool>(), false);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Splat_i32) {
@@ -220,12 +293,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::I32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 99);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, 99);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, 99);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AInt>(), 99);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AInt>(), 99);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AInt>(), 99);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Splat_u32) {
@@ -240,12 +326,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::U32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::U32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 99u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, 99u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, 99u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AInt>(), 99u);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AInt>(), 99u);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AInt>(), 99u);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Splat_f32) {
@@ -260,12 +359,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 9.9f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 9.9f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 9.9f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AFloat>(), 9.9f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AFloat>(), 9.9f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AFloat>(), 9.9f);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Splat_f16) {
@@ -281,13 +393,26 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F16>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
     // 9.9 is not exactly representable by f16, and should be quantized to 9.8984375
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 9.8984375f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 9.8984375f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 9.8984375f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AFloat>(), 9.8984375f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AFloat>(), 9.8984375f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AFloat>(), 9.8984375f);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Splat_bool) {
@@ -302,12 +427,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::Bool>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::Bool>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(0), true);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(1), true);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(2), true);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<bool>(), true);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<bool>(), true);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<bool>(), true);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_FullConstruct_i32) {
@@ -322,12 +460,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::I32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 1);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, 2);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, 3);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AInt>(), 1);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AInt>(), 2);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AInt>(), 3);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_FullConstruct_u32) {
@@ -342,12 +493,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::U32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::U32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 1);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, 2);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, 3);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AInt>(), 1);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AInt>(), 2);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AInt>(), 3);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_FullConstruct_f32) {
@@ -362,12 +526,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 1.f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 2.f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 3.f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AFloat>(), 1.f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AFloat>(), 2.f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AFloat>(), 3.f);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_FullConstruct_f16) {
@@ -383,12 +560,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F16>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 1.f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 2.f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 3.f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AFloat>(), 1.f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AFloat>(), 2.f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AFloat>(), 3.f);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_FullConstruct_bool) {
@@ -403,12 +593,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::Bool>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::Bool>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(0), true);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(1), false);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(2), true);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<bool>(), true);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<bool>(), false);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<bool>(), true);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_MixConstruct_i32) {
@@ -423,12 +626,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::I32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 1);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, 2);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, 3);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AInt>(), 1);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AInt>(), 2);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AInt>(), 3);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_MixConstruct_u32) {
@@ -443,12 +659,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::U32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::U32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 1);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, 2);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, 3);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AInt>(), 1);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AInt>(), 2);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AInt>(), 3);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_MixConstruct_f32) {
@@ -463,12 +692,157 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 1.f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 2.f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 3.f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AFloat>(), 1.f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AFloat>(), 2.f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AFloat>(), 3.f);
+}
+
+TEST_F(ResolverConstantsTest, Vec3_MixConstruct_f32_all_10) {
+    auto* expr = vec3<f32>(10_f, vec2<f32>(10_f, 10_f));
+    WrapInFunction(expr);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem = Sem().Get(expr);
+    EXPECT_NE(sem, nullptr);
+    auto* vec = sem->Type()->As<sem::Vector>();
+    ASSERT_NE(vec, nullptr);
+    EXPECT_TRUE(vec->type()->Is<sem::F32>());
+    EXPECT_EQ(vec->Width(), 3u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f32>(), 10_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f32>(), 10_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f32>(), 10_f);
+}
+
+TEST_F(ResolverConstantsTest, Vec3_MixConstruct_f32_all_positive_0) {
+    auto* expr = vec3<f32>(0_f, vec2<f32>(0_f, 0_f));
+    WrapInFunction(expr);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem = Sem().Get(expr);
+    EXPECT_NE(sem, nullptr);
+    auto* vec = sem->Type()->As<sem::Vector>();
+    ASSERT_NE(vec, nullptr);
+    EXPECT_TRUE(vec->type()->Is<sem::F32>());
+    EXPECT_EQ(vec->Width(), 3u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f32>(), 0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f32>(), 0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f32>(), 0_f);
+}
+
+TEST_F(ResolverConstantsTest, Vec3_MixConstruct_f32_all_negative_0) {
+    auto* expr = vec3<f32>(vec2<f32>(-0_f, -0_f), -0_f);
+    WrapInFunction(expr);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem = Sem().Get(expr);
+    EXPECT_NE(sem, nullptr);
+    auto* vec = sem->Type()->As<sem::Vector>();
+    ASSERT_NE(vec, nullptr);
+    EXPECT_TRUE(vec->type()->Is<sem::F32>());
+    EXPECT_EQ(vec->Width(), 3u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f32>(), -0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f32>(), -0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f32>(), -0_f);
+}
+
+TEST_F(ResolverConstantsTest, Vec3_MixConstruct_f32_mixed_sign_0) {
+    auto* expr = vec3<f32>(0_f, vec2<f32>(-0_f, 0_f));
+    WrapInFunction(expr);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem = Sem().Get(expr);
+    EXPECT_NE(sem, nullptr);
+    auto* vec = sem->Type()->As<sem::Vector>();
+    ASSERT_NE(vec, nullptr);
+    EXPECT_TRUE(vec->type()->Is<sem::F32>());
+    EXPECT_EQ(vec->Width(), 3u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f32>(), 0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f32>(), -0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f32>(), 0_f);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_MixConstruct_f16) {
@@ -484,12 +858,161 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F16>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 1.f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 2.f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 3.f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AFloat>(), 1.f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AFloat>(), 2.f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AFloat>(), 3.f);
+}
+
+TEST_F(ResolverConstantsTest, Vec3_MixConstruct_f16_all_10) {
+    Enable(ast::Extension::kF16);
+    auto* expr = vec3<f16>(10_h, vec2<f16>(10_h, 10_h));
+    WrapInFunction(expr);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem = Sem().Get(expr);
+    EXPECT_NE(sem, nullptr);
+    auto* vec = sem->Type()->As<sem::Vector>();
+    ASSERT_NE(vec, nullptr);
+    EXPECT_TRUE(vec->type()->Is<sem::F16>());
+    EXPECT_EQ(vec->Width(), 3u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f16>(), 10_h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f16>(), 10_h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f16>(), 10_h);
+}
+
+TEST_F(ResolverConstantsTest, Vec3_MixConstruct_f16_all_positive_0) {
+    Enable(ast::Extension::kF16);
+    auto* expr = vec3<f16>(0_h, vec2<f16>(0_h, 0_h));
+    WrapInFunction(expr);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem = Sem().Get(expr);
+    EXPECT_NE(sem, nullptr);
+    auto* vec = sem->Type()->As<sem::Vector>();
+    ASSERT_NE(vec, nullptr);
+    EXPECT_TRUE(vec->type()->Is<sem::F16>());
+    EXPECT_EQ(vec->Width(), 3u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f16>(), 0_h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f16>(), 0_h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f16>(), 0_h);
+}
+
+TEST_F(ResolverConstantsTest, Vec3_MixConstruct_f16_all_negative_0) {
+    Enable(ast::Extension::kF16);
+    auto* expr = vec3<f16>(vec2<f16>(-0_h, -0_h), -0_h);
+    WrapInFunction(expr);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem = Sem().Get(expr);
+    EXPECT_NE(sem, nullptr);
+    auto* vec = sem->Type()->As<sem::Vector>();
+    ASSERT_NE(vec, nullptr);
+    EXPECT_TRUE(vec->type()->Is<sem::F16>());
+    EXPECT_EQ(vec->Width(), 3u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f16>(), -0_h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f16>(), -0_h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f16>(), -0_h);
+}
+
+TEST_F(ResolverConstantsTest, Vec3_MixConstruct_f16_mixed_sign_0) {
+    Enable(ast::Extension::kF16);
+    auto* expr = vec3<f16>(0_h, vec2<f16>(-0_h, 0_h));
+    WrapInFunction(expr);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem = Sem().Get(expr);
+    EXPECT_NE(sem, nullptr);
+    auto* vec = sem->Type()->As<sem::Vector>();
+    ASSERT_NE(vec, nullptr);
+    EXPECT_TRUE(vec->type()->Is<sem::F16>());
+    EXPECT_EQ(vec->Width(), 3u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f16>(), 0_h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f16>(), -0_h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f16>(), 0_h);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_MixConstruct_bool) {
@@ -504,12 +1027,91 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::Bool>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::Bool>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(0), true);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(1), false);
-    EXPECT_EQ(sem->ConstantValue().Element<bool>(2), true);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<bool>(), true);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<bool>(), false);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<bool>(), true);
+}
+
+TEST_F(ResolverConstantsTest, Vec3_MixConstruct_all_true) {
+    auto* expr = vec3<bool>(true, vec2<bool>(true, true));
+    WrapInFunction(expr);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem = Sem().Get(expr);
+    ASSERT_NE(sem, nullptr);
+    auto* vec = sem->Type()->As<sem::Vector>();
+    ASSERT_NE(vec, nullptr);
+    EXPECT_TRUE(vec->type()->Is<sem::Bool>());
+    EXPECT_EQ(vec->Width(), 3u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<bool>(), true);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<bool>(), true);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<bool>(), true);
+}
+
+TEST_F(ResolverConstantsTest, Vec3_MixConstruct_all_false) {
+    auto* expr = vec3<bool>(false, vec2<bool>(false, false));
+    WrapInFunction(expr);
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+    auto* sem = Sem().Get(expr);
+    ASSERT_NE(sem, nullptr);
+    auto* vec = sem->Type()->As<sem::Vector>();
+    ASSERT_NE(vec, nullptr);
+    EXPECT_TRUE(vec->type()->Is<sem::Bool>());
+    EXPECT_EQ(vec->Width(), 3u);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<bool>(), false);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<bool>(), false);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<bool>(), false);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Convert_f32_to_i32) {
@@ -524,12 +1126,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::I32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 1);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, 2);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, 3);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AInt>(), 1);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AInt>(), 2);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AInt>(), 3);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Convert_u32_to_f32) {
@@ -544,12 +1159,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 10.f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 20.f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 30.f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AFloat>(), 10.f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AFloat>(), 20.f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AFloat>(), 30.f);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Convert_f16_to_i32) {
@@ -565,12 +1193,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::I32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, 1);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, 2);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, 3);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AInt>(), 1_i);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AInt>(), 2_i);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AInt>(), 3_i);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Convert_u32_to_f16) {
@@ -586,12 +1227,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F16>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 10.f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 20.f);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 30.f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AFloat>(), 10.f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AFloat>(), 20.f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AFloat>(), 30.f);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Convert_Large_f32_to_i32) {
@@ -606,12 +1260,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::I32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, i32::kHighest);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, i32::kLowest);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, i32::kHighest);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AInt>(), i32::kHighest);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AInt>(), i32::kLowest);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AInt>(), i32::kHighest);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Convert_Large_f32_to_u32) {
@@ -626,12 +1293,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::U32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::U32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(0).value, u32::kHighest);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(1).value, u32::kLowest);
-    EXPECT_EQ(sem->ConstantValue().Element<AInt>(2).value, u32::kHighest);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AInt>(), u32::kHighest);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AInt>(), u32::kLowest);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AInt>(), u32::kHighest);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Convert_Large_f32_to_f16) {
@@ -650,12 +1330,25 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F16>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, kInf);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, -kInf);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, kInf);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AFloat>(), kInf);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AFloat>(), -kInf);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AFloat>(), kInf);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Convert_Small_f32_to_f16) {
@@ -672,12 +1365,28 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F16>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 0.0);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, -0.0);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 0.0);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<AFloat>(), 0.0);
+    EXPECT_FALSE(std::signbit(sem->ConstantValue()->Index(0)->As<AFloat>().value));
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<AFloat>(), -0.0);
+    EXPECT_TRUE(std::signbit(sem->ConstantValue()->Index(1)->As<AFloat>().value));
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<AFloat>(), 0.0);
+    EXPECT_FALSE(std::signbit(sem->ConstantValue()->Index(2)->As<AFloat>().value));
 }
 
 TEST_F(ResolverConstantsTest, Mat2x3_ZeroInit_f32) {
@@ -693,15 +1402,40 @@
     EXPECT_TRUE(mat->type()->Is<sem::F32>());
     EXPECT_EQ(mat->columns(), 2u);
     EXPECT_EQ(mat->rows(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 6u);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(0).value, 0._f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(1).value, 0._f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(2).value, 0._f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(3).value, 0._f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(4).value, 0._f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(5).value, 0._f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(0)->As<f32>(), 0._f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(1)->As<f32>(), 0._f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(2)->As<f32>(), 0._f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(0)->As<f32>(), 0._f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(1)->As<f32>(), 0._f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(2)->As<f32>(), 0._f);
 }
 
 TEST_F(ResolverConstantsTest, Mat2x3_ZeroInit_f16) {
@@ -719,15 +1453,40 @@
     EXPECT_TRUE(mat->type()->Is<sem::F16>());
     EXPECT_EQ(mat->columns(), 2u);
     EXPECT_EQ(mat->rows(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F16>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 6u);
-    EXPECT_EQ(sem->ConstantValue().Element<f16>(0).value, 0._h);
-    EXPECT_EQ(sem->ConstantValue().Element<f16>(1).value, 0._h);
-    EXPECT_EQ(sem->ConstantValue().Element<f16>(2).value, 0._h);
-    EXPECT_EQ(sem->ConstantValue().Element<f16>(3).value, 0._h);
-    EXPECT_EQ(sem->ConstantValue().Element<f16>(4).value, 0._h);
-    EXPECT_EQ(sem->ConstantValue().Element<f16>(5).value, 0._h);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(0)->As<f16>(), 0._h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(1)->As<f16>(), 0._h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(2)->As<f16>(), 0._h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(0)->As<f16>(), 0._h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(1)->As<f16>(), 0._h);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(2)->As<f16>(), 0._h);
 }
 
 TEST_F(ResolverConstantsTest, Mat3x2_Construct_Scalars_af) {
@@ -743,15 +1502,40 @@
     EXPECT_TRUE(mat->type()->Is<sem::F32>());
     EXPECT_EQ(mat->columns(), 3u);
     EXPECT_EQ(mat->rows(), 2u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 6u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 1._a);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 2._a);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 3._a);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(3).value, 4._a);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(4).value, 5._a);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(5).value, 6._a);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(0)->As<AFloat>(), 1._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(1)->As<AFloat>(), 2._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(0)->As<AFloat>(), 3._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(1)->As<AFloat>(), 4._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->Index(0)->As<AFloat>(), 5._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->Index(1)->As<AFloat>(), 6._a);
 }
 
 TEST_F(ResolverConstantsTest, Mat3x2_Construct_Columns_af) {
@@ -770,15 +1554,40 @@
     EXPECT_TRUE(mat->type()->Is<sem::F32>());
     EXPECT_EQ(mat->columns(), 3u);
     EXPECT_EQ(mat->rows(), 2u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 6u);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(0).value, 1._a);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(1).value, 2._a);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(2).value, 3._a);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(3).value, 4._a);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(4).value, 5._a);
-    EXPECT_EQ(sem->ConstantValue().Element<AFloat>(5).value, 6._a);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(0)->As<AFloat>(), 1._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(1)->As<AFloat>(), 2._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(0)->As<AFloat>(), 3._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(1)->As<AFloat>(), 4._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->Index(0)->As<AFloat>(), 5._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->Index(1)->As<AFloat>(), 6._a);
 }
 
 TEST_F(ResolverConstantsTest, Array_i32_Zero) {
@@ -793,13 +1602,30 @@
     ASSERT_NE(arr, nullptr);
     EXPECT_TRUE(arr->ElemType()->Is<sem::I32>());
     EXPECT_EQ(arr->Count(), 4u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 4u);
-    EXPECT_EQ(sem->ConstantValue().Element<i32>(0).value, 0_i);
-    EXPECT_EQ(sem->ConstantValue().Element<i32>(1).value, 0_i);
-    EXPECT_EQ(sem->ConstantValue().Element<i32>(2).value, 0_i);
-    EXPECT_EQ(sem->ConstantValue().Element<i32>(3).value, 0_i);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<i32>(), 0_i);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<i32>(), 0_i);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<i32>(), 0_i);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(3)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(3)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(3)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(3)->As<i32>(), 0_i);
 }
 
 TEST_F(ResolverConstantsTest, Array_f32_Zero) {
@@ -814,13 +1640,30 @@
     ASSERT_NE(arr, nullptr);
     EXPECT_TRUE(arr->ElemType()->Is<sem::F32>());
     EXPECT_EQ(arr->Count(), 4u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 4u);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(0).value, 0_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(1).value, 0_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(2).value, 0_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(3).value, 0_f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f32>(), 0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f32>(), 0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f32>(), 0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(3)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(3)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(3)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(3)->As<f32>(), 0_f);
 }
 
 TEST_F(ResolverConstantsTest, Array_vec3_f32_Zero) {
@@ -835,15 +1678,40 @@
     ASSERT_NE(arr, nullptr);
     EXPECT_TRUE(arr->ElemType()->Is<sem::Vector>());
     EXPECT_EQ(arr->Count(), 2u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 6u);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(0).value, 0_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(1).value, 0_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(2).value, 0_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(3).value, 0_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(4).value, 0_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(5).value, 0_f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(0)->As<f32>(), 0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(1)->As<f32>(), 0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(2)->As<f32>(), 0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(0)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(0)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(0)->As<f32>(), 0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(1)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(1)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(1)->As<f32>(), 0_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(2)->AllEqual());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(2)->AnyZero());
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(2)->As<f32>(), 0_f);
 }
 
 TEST_F(ResolverConstantsTest, Array_i32_Elements) {
@@ -858,13 +1726,30 @@
     ASSERT_NE(arr, nullptr);
     EXPECT_TRUE(arr->ElemType()->Is<sem::I32>());
     EXPECT_EQ(arr->Count(), 4u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 4u);
-    EXPECT_EQ(sem->ConstantValue().Element<i32>(0).value, 10_i);
-    EXPECT_EQ(sem->ConstantValue().Element<i32>(1).value, 20_i);
-    EXPECT_EQ(sem->ConstantValue().Element<i32>(2).value, 30_i);
-    EXPECT_EQ(sem->ConstantValue().Element<i32>(3).value, 40_i);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<i32>(), 10_i);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<i32>(), 20_i);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<i32>(), 30_i);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(3)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(3)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(3)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(3)->As<i32>(), 40_i);
 }
 
 TEST_F(ResolverConstantsTest, Array_f32_Elements) {
@@ -879,13 +1764,30 @@
     ASSERT_NE(arr, nullptr);
     EXPECT_TRUE(arr->ElemType()->Is<sem::F32>());
     EXPECT_EQ(arr->Count(), 4u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 4u);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(0).value, 10_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(1).value, 20_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(2).value, 30_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(3).value, 40_f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f32>(), 10_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f32>(), 20_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f32>(), 30_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(3)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(3)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(3)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(3)->As<f32>(), 40_f);
 }
 
 TEST_F(ResolverConstantsTest, Array_vec3_f32_Elements) {
@@ -901,15 +1803,16 @@
     ASSERT_NE(arr, nullptr);
     EXPECT_TRUE(arr->ElemType()->Is<sem::Vector>());
     EXPECT_EQ(arr->Count(), 2u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 6u);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(0).value, 1_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(1).value, 2_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(2).value, 3_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(3).value, 4_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(4).value, 5_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(5).value, 6_f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(0)->As<f32>(), 1_f);
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(1)->As<f32>(), 2_f);
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->Index(2)->As<f32>(), 3_f);
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(0)->As<f32>(), 4_f);
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(1)->As<f32>(), 5_f);
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->Index(2)->As<f32>(), 6_f);
 }
 
 ////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -925,10 +1828,11 @@
     auto* sem = Sem().Get(expr);
     ASSERT_NE(sem, nullptr);
     ASSERT_TRUE(sem->Type()->Is<sem::I32>());
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 1u);
-    EXPECT_EQ(sem->ConstantValue().Element<i32>(0).value, 3_i);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->As<i32>(), 3_i);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Index_OOB_High) {
@@ -941,10 +1845,11 @@
     auto* sem = Sem().Get(expr);
     ASSERT_NE(sem, nullptr);
     ASSERT_TRUE(sem->Type()->Is<sem::I32>());
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 1u);
-    EXPECT_EQ(sem->ConstantValue().Element<i32>(0).value, 3_i);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->As<i32>(), 3_i);
 }
 
 TEST_F(ResolverConstantsTest, Vec3_Index_OOB_Low) {
@@ -957,10 +1862,11 @@
     auto* sem = Sem().Get(expr);
     ASSERT_NE(sem, nullptr);
     ASSERT_TRUE(sem->Type()->Is<sem::I32>());
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::I32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 1u);
-    EXPECT_EQ(sem->ConstantValue().Element<i32>(0).value, 1_i);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_TRUE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->As<i32>(), 1_i);
 }
 
 TEST_F(ResolverConstantsTest, Mat3x2_Index) {
@@ -975,11 +1881,17 @@
     auto* vec = sem->Type()->As<sem::Vector>();
     ASSERT_NE(vec, nullptr);
     EXPECT_EQ(vec->Width(), 2u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 2u);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(0).value, 5._a);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(1).value, 6._a);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f32>(), 5._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f32>(), 6._a);
 }
 
 TEST_F(ResolverConstantsTest, Mat3x2_Index_OOB_High) {
@@ -996,11 +1908,17 @@
     auto* vec = sem->Type()->As<sem::Vector>();
     ASSERT_NE(vec, nullptr);
     EXPECT_EQ(vec->Width(), 2u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 2u);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(0).value, 5._a);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(1).value, 6._a);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f32>(), 5._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f32>(), 6._a);
 }
 
 TEST_F(ResolverConstantsTest, Mat3x2_Index_OOB_Low) {
@@ -1017,11 +1935,17 @@
     auto* vec = sem->Type()->As<sem::Vector>();
     ASSERT_NE(vec, nullptr);
     EXPECT_EQ(vec->Width(), 2u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 2u);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(0).value, 1._a);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(1).value, 2._a);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f32>(), 1._a);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f32>(), 2._a);
 }
 
 TEST_F(ResolverConstantsTest, Array_vec3_f32_Index) {
@@ -1038,12 +1962,22 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(0).value, 4_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(1).value, 5_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(2).value, 6_f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f32>(), 4_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f32>(), 5_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f32>(), 6_f);
 }
 
 TEST_F(ResolverConstantsTest, Array_vec3_f32_Index_OOB_High) {
@@ -1061,12 +1995,22 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(0).value, 4_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(1).value, 5_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(2).value, 6_f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f32>(), 4_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f32>(), 5_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f32>(), 6_f);
 }
 
 TEST_F(ResolverConstantsTest, Array_vec3_f32_Index_OOB_Low) {
@@ -1084,19 +2028,32 @@
     ASSERT_NE(vec, nullptr);
     EXPECT_TRUE(vec->type()->Is<sem::F32>());
     EXPECT_EQ(vec->Width(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Type(), sem->Type());
-    EXPECT_TRUE(sem->ConstantValue().ElementType()->Is<sem::F32>());
-    ASSERT_EQ(sem->ConstantValue().ElementCount(), 3u);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(0).value, 1_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(1).value, 2_f);
-    EXPECT_EQ(sem->ConstantValue().Element<f32>(2).value, 3_f);
+    EXPECT_TYPE(sem->ConstantValue()->Type(), sem->Type());
+    EXPECT_FALSE(sem->ConstantValue()->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->AllZero());
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(0)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(0)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(0)->As<f32>(), 1_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(1)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(1)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(1)->As<f32>(), 2_f);
+
+    EXPECT_TRUE(sem->ConstantValue()->Index(2)->AllEqual());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AnyZero());
+    EXPECT_FALSE(sem->ConstantValue()->Index(2)->AllZero());
+    EXPECT_EQ(sem->ConstantValue()->Index(2)->As<f32>(), 3_f);
 }
 
 TEST_F(ResolverConstantsTest, ChainedIndex) {
     auto* arr_expr = Construct(ty.array(ty.mat2x3<f32>(), 2_u),        // array<mat2x3<f32>, 2u>
                                mat2x3<f32>(vec3<f32>(1_f, 2_f, 3_f),   //
                                            vec3<f32>(4_f, 5_f, 6_f)),  //
-                               mat2x3<f32>(vec3<f32>(7_f, 8_f, 9_f),   //
+                               mat2x3<f32>(vec3<f32>(7_f, 0_f, 9_f),   //
                                            vec3<f32>(10_f, 11_f, 12_f)));
 
     auto* mat_expr = IndexAccessor(arr_expr, 1_i);  // arr[1]
@@ -1114,15 +2071,40 @@
         EXPECT_TRUE(ty->ColumnType()->Is<sem::Vector>());
         EXPECT_EQ(ty->columns(), 2u);
         EXPECT_EQ(ty->rows(), 3u);
-        EXPECT_EQ(mat->ConstantValue().Type(), mat->Type());
-        EXPECT_TRUE(mat->ConstantValue().ElementType()->Is<sem::F32>());
-        ASSERT_EQ(mat->ConstantValue().ElementCount(), 6u);
-        EXPECT_EQ(mat->ConstantValue().Element<f32>(0).value, 7_f);
-        EXPECT_EQ(mat->ConstantValue().Element<f32>(1).value, 8_f);
-        EXPECT_EQ(mat->ConstantValue().Element<f32>(2).value, 9_f);
-        EXPECT_EQ(mat->ConstantValue().Element<f32>(3).value, 10_f);
-        EXPECT_EQ(mat->ConstantValue().Element<f32>(4).value, 11_f);
-        EXPECT_EQ(mat->ConstantValue().Element<f32>(5).value, 12_f);
+        EXPECT_EQ(mat->ConstantValue()->Type(), mat->Type());
+        EXPECT_FALSE(mat->ConstantValue()->AllEqual());
+        EXPECT_TRUE(mat->ConstantValue()->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->AllZero());
+
+        EXPECT_TRUE(mat->ConstantValue()->Index(0)->Index(0)->AllEqual());
+        EXPECT_FALSE(mat->ConstantValue()->Index(0)->Index(0)->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->Index(0)->Index(0)->AllZero());
+        EXPECT_EQ(mat->ConstantValue()->Index(0)->Index(0)->As<f32>(), 7_f);
+
+        EXPECT_TRUE(mat->ConstantValue()->Index(0)->Index(1)->AllEqual());
+        EXPECT_TRUE(mat->ConstantValue()->Index(0)->Index(1)->AnyZero());
+        EXPECT_TRUE(mat->ConstantValue()->Index(0)->Index(1)->AllZero());
+        EXPECT_EQ(mat->ConstantValue()->Index(0)->Index(1)->As<f32>(), 0_f);
+
+        EXPECT_TRUE(mat->ConstantValue()->Index(0)->Index(2)->AllEqual());
+        EXPECT_FALSE(mat->ConstantValue()->Index(0)->Index(2)->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->Index(0)->Index(2)->AllZero());
+        EXPECT_EQ(mat->ConstantValue()->Index(0)->Index(2)->As<f32>(), 9_f);
+
+        EXPECT_TRUE(mat->ConstantValue()->Index(1)->Index(0)->AllEqual());
+        EXPECT_FALSE(mat->ConstantValue()->Index(1)->Index(0)->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->Index(1)->Index(0)->AllZero());
+        EXPECT_EQ(mat->ConstantValue()->Index(1)->Index(0)->As<f32>(), 10_f);
+
+        EXPECT_TRUE(mat->ConstantValue()->Index(1)->Index(1)->AllEqual());
+        EXPECT_FALSE(mat->ConstantValue()->Index(1)->Index(1)->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->Index(1)->Index(1)->AllZero());
+        EXPECT_EQ(mat->ConstantValue()->Index(1)->Index(1)->As<f32>(), 11_f);
+
+        EXPECT_TRUE(mat->ConstantValue()->Index(1)->Index(2)->AllEqual());
+        EXPECT_FALSE(mat->ConstantValue()->Index(1)->Index(2)->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->Index(1)->Index(2)->AllZero());
+        EXPECT_EQ(mat->ConstantValue()->Index(1)->Index(2)->As<f32>(), 12_f);
     }
     {
         auto* vec = Sem().Get(vec_expr);
@@ -1131,21 +2113,35 @@
         ASSERT_NE(vec->Type(), nullptr);
         EXPECT_TRUE(ty->type()->Is<sem::F32>());
         EXPECT_EQ(ty->Width(), 3u);
-        EXPECT_EQ(vec->ConstantValue().Type(), vec->Type());
-        EXPECT_TRUE(vec->ConstantValue().ElementType()->Is<sem::F32>());
-        ASSERT_EQ(vec->ConstantValue().ElementCount(), 3u);
-        EXPECT_EQ(vec->ConstantValue().Element<f32>(0).value, 7_f);
-        EXPECT_EQ(vec->ConstantValue().Element<f32>(1).value, 8_f);
-        EXPECT_EQ(vec->ConstantValue().Element<f32>(2).value, 9_f);
+        EXPECT_EQ(vec->ConstantValue()->Type(), vec->Type());
+        EXPECT_FALSE(vec->ConstantValue()->AllEqual());
+        EXPECT_TRUE(vec->ConstantValue()->AnyZero());
+        EXPECT_FALSE(vec->ConstantValue()->AllZero());
+
+        EXPECT_TRUE(vec->ConstantValue()->Index(0)->AllEqual());
+        EXPECT_FALSE(vec->ConstantValue()->Index(0)->AnyZero());
+        EXPECT_FALSE(vec->ConstantValue()->Index(0)->AllZero());
+        EXPECT_EQ(vec->ConstantValue()->Index(0)->As<f32>(), 7_f);
+
+        EXPECT_TRUE(vec->ConstantValue()->Index(1)->AllEqual());
+        EXPECT_TRUE(vec->ConstantValue()->Index(1)->AnyZero());
+        EXPECT_TRUE(vec->ConstantValue()->Index(1)->AllZero());
+        EXPECT_EQ(vec->ConstantValue()->Index(1)->As<f32>(), 0_f);
+
+        EXPECT_TRUE(vec->ConstantValue()->Index(2)->AllEqual());
+        EXPECT_FALSE(vec->ConstantValue()->Index(2)->AnyZero());
+        EXPECT_FALSE(vec->ConstantValue()->Index(2)->AllZero());
+        EXPECT_EQ(vec->ConstantValue()->Index(2)->As<f32>(), 9_f);
     }
     {
         auto* f = Sem().Get(f32_expr);
         EXPECT_NE(f, nullptr);
         EXPECT_TRUE(f->Type()->Is<sem::F32>());
-        EXPECT_EQ(f->ConstantValue().Type(), f->Type());
-        EXPECT_TRUE(f->ConstantValue().ElementType()->Is<sem::F32>());
-        ASSERT_EQ(f->ConstantValue().ElementCount(), 1u);
-        EXPECT_EQ(f->ConstantValue().Element<f32>(0).value, 9_f);
+        EXPECT_EQ(f->ConstantValue()->Type(), f->Type());
+        EXPECT_TRUE(f->ConstantValue()->AllEqual());
+        EXPECT_FALSE(f->ConstantValue()->AnyZero());
+        EXPECT_FALSE(f->ConstantValue()->AllZero());
+        EXPECT_EQ(f->ConstantValue()->As<f32>(), 9_f);
     }
 }
 
@@ -1174,15 +2170,40 @@
         EXPECT_TRUE(ty->ColumnType()->Is<sem::Vector>());
         EXPECT_EQ(ty->columns(), 2u);
         EXPECT_EQ(ty->rows(), 3u);
-        EXPECT_EQ(mat->ConstantValue().Type(), mat->Type());
-        EXPECT_TRUE(mat->ConstantValue().ElementType()->Is<sem::F32>());
-        ASSERT_EQ(mat->ConstantValue().ElementCount(), 6u);
-        EXPECT_EQ(mat->ConstantValue().Element<f32>(0).value, 1_f);
-        EXPECT_EQ(mat->ConstantValue().Element<f32>(1).value, 2_f);
-        EXPECT_EQ(mat->ConstantValue().Element<f32>(2).value, 3_f);
-        EXPECT_EQ(mat->ConstantValue().Element<f32>(3).value, 4_f);
-        EXPECT_EQ(mat->ConstantValue().Element<f32>(4).value, 5_f);
-        EXPECT_EQ(mat->ConstantValue().Element<f32>(5).value, 6_f);
+        EXPECT_EQ(mat->ConstantValue()->Type(), mat->Type());
+        EXPECT_FALSE(mat->ConstantValue()->AllEqual());
+        EXPECT_FALSE(mat->ConstantValue()->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->AllZero());
+
+        EXPECT_TRUE(mat->ConstantValue()->Index(0)->Index(0)->AllEqual());
+        EXPECT_FALSE(mat->ConstantValue()->Index(0)->Index(0)->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->Index(0)->Index(0)->AllZero());
+        EXPECT_EQ(mat->ConstantValue()->Index(0)->Index(0)->As<f32>(), 1_f);
+
+        EXPECT_TRUE(mat->ConstantValue()->Index(0)->Index(1)->AllEqual());
+        EXPECT_FALSE(mat->ConstantValue()->Index(0)->Index(1)->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->Index(0)->Index(1)->AllZero());
+        EXPECT_EQ(mat->ConstantValue()->Index(0)->Index(1)->As<f32>(), 2_f);
+
+        EXPECT_TRUE(mat->ConstantValue()->Index(0)->Index(2)->AllEqual());
+        EXPECT_FALSE(mat->ConstantValue()->Index(0)->Index(2)->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->Index(0)->Index(2)->AllZero());
+        EXPECT_EQ(mat->ConstantValue()->Index(0)->Index(2)->As<f32>(), 3_f);
+
+        EXPECT_TRUE(mat->ConstantValue()->Index(1)->Index(0)->AllEqual());
+        EXPECT_FALSE(mat->ConstantValue()->Index(1)->Index(0)->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->Index(1)->Index(0)->AllZero());
+        EXPECT_EQ(mat->ConstantValue()->Index(1)->Index(0)->As<f32>(), 4_f);
+
+        EXPECT_TRUE(mat->ConstantValue()->Index(1)->Index(1)->AllEqual());
+        EXPECT_FALSE(mat->ConstantValue()->Index(1)->Index(1)->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->Index(1)->Index(1)->AllZero());
+        EXPECT_EQ(mat->ConstantValue()->Index(1)->Index(1)->As<f32>(), 5_f);
+
+        EXPECT_TRUE(mat->ConstantValue()->Index(1)->Index(2)->AllEqual());
+        EXPECT_FALSE(mat->ConstantValue()->Index(1)->Index(2)->AnyZero());
+        EXPECT_FALSE(mat->ConstantValue()->Index(1)->Index(2)->AllZero());
+        EXPECT_EQ(mat->ConstantValue()->Index(1)->Index(2)->As<f32>(), 6_f);
     }
     {
         auto* vec = Sem().Get(vec_expr);
@@ -1191,21 +2212,35 @@
         ASSERT_NE(vec->Type(), nullptr);
         EXPECT_TRUE(ty->type()->Is<sem::F32>());
         EXPECT_EQ(ty->Width(), 3u);
-        EXPECT_EQ(vec->ConstantValue().Type(), vec->Type());
-        EXPECT_TRUE(vec->ConstantValue().ElementType()->Is<sem::F32>());
-        ASSERT_EQ(vec->ConstantValue().ElementCount(), 3u);
-        EXPECT_EQ(vec->ConstantValue().Element<f32>(0).value, 1_f);
-        EXPECT_EQ(vec->ConstantValue().Element<f32>(1).value, 2_f);
-        EXPECT_EQ(vec->ConstantValue().Element<f32>(2).value, 3_f);
+        EXPECT_EQ(vec->ConstantValue()->Type(), vec->Type());
+        EXPECT_FALSE(vec->ConstantValue()->AllEqual());
+        EXPECT_FALSE(vec->ConstantValue()->AnyZero());
+        EXPECT_FALSE(vec->ConstantValue()->AllZero());
+
+        EXPECT_TRUE(vec->ConstantValue()->Index(0)->AllEqual());
+        EXPECT_FALSE(vec->ConstantValue()->Index(0)->AnyZero());
+        EXPECT_FALSE(vec->ConstantValue()->Index(0)->AllZero());
+        EXPECT_EQ(vec->ConstantValue()->Index(0)->As<f32>(), 1_f);
+
+        EXPECT_TRUE(vec->ConstantValue()->Index(1)->AllEqual());
+        EXPECT_FALSE(vec->ConstantValue()->Index(1)->AnyZero());
+        EXPECT_FALSE(vec->ConstantValue()->Index(1)->AllZero());
+        EXPECT_EQ(vec->ConstantValue()->Index(1)->As<f32>(), 2_f);
+
+        EXPECT_TRUE(vec->ConstantValue()->Index(2)->AllEqual());
+        EXPECT_FALSE(vec->ConstantValue()->Index(2)->AnyZero());
+        EXPECT_FALSE(vec->ConstantValue()->Index(2)->AllZero());
+        EXPECT_EQ(vec->ConstantValue()->Index(2)->As<f32>(), 3_f);
     }
     {
         auto* f = Sem().Get(f32_expr);
         EXPECT_NE(f, nullptr);
         EXPECT_TRUE(f->Type()->Is<sem::F32>());
-        EXPECT_EQ(f->ConstantValue().Type(), f->Type());
-        EXPECT_TRUE(f->ConstantValue().ElementType()->Is<sem::F32>());
-        ASSERT_EQ(f->ConstantValue().ElementCount(), 1u);
-        EXPECT_EQ(f->ConstantValue().Element<f32>(0).value, 3_f);
+        EXPECT_EQ(f->ConstantValue()->Type(), f->Type());
+        EXPECT_TRUE(f->ConstantValue()->AllEqual());
+        EXPECT_FALSE(f->ConstantValue()->AnyZero());
+        EXPECT_FALSE(f->ConstantValue()->AllZero());
+        EXPECT_EQ(f->ConstantValue()->As<f32>(), 3_f);
     }
 }
 
diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc
index 0048d74..558cb12 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -1604,16 +1604,16 @@
     auto& signature = builtin->Signature();
 
     auto check_arg_is_constexpr = [&](sem::ParameterUsage usage, int min, int max) {
-        auto index = signature.IndexOf(usage);
-        if (index < 0) {
+        auto signed_index = signature.IndexOf(usage);
+        if (signed_index < 0) {
             return true;
         }
+        auto index = static_cast<size_t>(signed_index);
         std::string name = sem::str(usage);
-        auto* arg = call->Arguments()[static_cast<size_t>(index)];
+        auto* arg = call->Arguments()[index];
         if (auto values = arg->ConstantValue()) {
             // Assert that the constant values are of the expected type.
-            if (!values.Type()->IsAnyOf<sem::I32, sem::Vector>() ||
-                !values.ElementType()->Is<sem::I32>()) {
+            if (!values->Type()->is_integer_scalar_or_vector()) {
                 TINT_ICE(Resolver, diagnostics_)
                     << "failed to resolve '" + func_name + "' " << name << " parameter type";
                 return false;
@@ -1631,25 +1631,26 @@
                     return ast::TraverseAction::Stop;
                 });
             if (is_const_expr) {
-                auto vector =
-                    builtin->Parameters()[static_cast<size_t>(index)]->Type()->Is<sem::Vector>();
-                for (size_t i = 0, n = values.ElementCount(); i < n; i++) {
-                    auto value = values.Element<AInt>(i).value;
-                    if (value < min || value > max) {
-                        if (vector) {
+                if (auto* vector = builtin->Parameters()[index]->Type()->As<sem::Vector>()) {
+                    for (size_t i = 0; i < vector->Width(); i++) {
+                        auto value = values->Index(i)->As<AInt>();
+                        if (value < min || value > max) {
                             AddError("each component of the " + name +
                                          " argument must be at least " + std::to_string(min) +
                                          " and at most " + std::to_string(max) + ". " + name +
                                          " component " + std::to_string(i) + " is " +
                                          std::to_string(value),
                                      arg->Declaration()->source);
-                        } else {
-                            AddError("the " + name + " argument must be at least " +
-                                         std::to_string(min) + " and at most " +
-                                         std::to_string(max) + ". " + name + " is " +
-                                         std::to_string(value),
-                                     arg->Declaration()->source);
+                            return false;
                         }
+                    }
+                } else {
+                    auto value = values->As<AInt>();
+                    if (value < min || value > max) {
+                        AddError("the " + name + " argument must be at least " +
+                                     std::to_string(min) + " and at most " + std::to_string(max) +
+                                     ". " + name + " is " + std::to_string(value),
+                                 arg->Declaration()->source);
                         return false;
                     }
                 }
diff --git a/src/tint/resolver/variable_test.cc b/src/tint/resolver/variable_test.cc
index 2ee5979..7eb0dea 100644
--- a/src/tint/resolver/variable_test.cc
+++ b/src/tint/resolver/variable_test.cc
@@ -920,21 +920,13 @@
     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);
+    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());
 }
 
 TEST_F(ResolverVariableTest, LocalConst_ImplicitType_Decls) {
@@ -949,7 +941,11 @@
     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)));
+    auto* c_maf32 = Const("l", nullptr,
+                          Construct(ty.mat(nullptr, 3, 3),  //
+                                    Construct(ty.vec(nullptr, 3), Expr(0._a)),
+                                    Construct(ty.vec(nullptr, 3), Expr(0._a)),
+                                    Construct(ty.vec(nullptr, 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);
@@ -982,31 +978,18 @@
     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);
+    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());
 }
 
 TEST_F(ResolverVariableTest, LocalConst_PropagateConstValue) {
@@ -1020,8 +1003,7 @@
 
     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);
+    EXPECT_EQ(Sem().Get(c)->ConstantValue()->As<i32>(), 42_i);
 }
 
 // Enable when we have @const operators implemented
@@ -1034,8 +1016,7 @@
 
     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);
+    EXPECT_EQ(Sem().Get(c)->ConstantValue()->As<i32>(), 3_i);
 }
 
 ////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -1126,21 +1107,13 @@
     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);
+    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());
 }
 
 TEST_F(ResolverVariableTest, GlobalConst_ImplicitType_Decls) {
@@ -1155,7 +1128,11 @@
     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)));
+    auto* c_maf32 = GlobalConst("l", nullptr,
+                                Construct(ty.mat(nullptr, 3, 3),  //
+                                          Construct(ty.vec(nullptr, 3), Expr(0._a)),
+                                          Construct(ty.vec(nullptr, 3), Expr(0._a)),
+                                          Construct(ty.vec(nullptr, 3), Expr(0._a))));
 
     ASSERT_TRUE(r()->Resolve()) << r()->error();
 
@@ -1185,31 +1162,18 @@
     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);
+    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());
 }
 
 TEST_F(ResolverVariableTest, GlobalConst_PropagateConstValue) {
@@ -1221,8 +1185,7 @@
 
     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);
+    EXPECT_EQ(Sem().Get(c)->ConstantValue()->As<i32>(), 42_i);
 }
 
 // Enable when we have @const operators implemented
@@ -1233,8 +1196,7 @@
 
     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);
+    EXPECT_EQ(Sem().Get(c)->ConstantValue()->As<i32>(), 3_i);
 }
 
 ////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/sem/call.cc b/src/tint/sem/call.cc
index f688cce..bfce3b1 100644
--- a/src/tint/sem/call.cc
+++ b/src/tint/sem/call.cc
@@ -25,9 +25,9 @@
            const CallTarget* target,
            std::vector<const sem::Expression*> arguments,
            const Statement* statement,
-           Constant constant,
+           const Constant* constant,
            bool has_side_effects)
-    : Base(declaration, target->ReturnType(), statement, std::move(constant), has_side_effects),
+    : Base(declaration, target->ReturnType(), statement, constant, has_side_effects),
       target_(target),
       arguments_(std::move(arguments)) {}
 
diff --git a/src/tint/sem/call.h b/src/tint/sem/call.h
index e27f388..1955bf9 100644
--- a/src/tint/sem/call.h
+++ b/src/tint/sem/call.h
@@ -38,7 +38,7 @@
          const CallTarget* target,
          std::vector<const sem::Expression*> arguments,
          const Statement* statement,
-         Constant constant,
+         const Constant* constant,
          bool has_side_effects);
 
     /// Destructor
diff --git a/src/tint/sem/constant.cc b/src/tint/sem/constant.cc
index ed38686..70bb08c 100644
--- a/src/tint/sem/constant.cc
+++ b/src/tint/sem/constant.cc
@@ -14,102 +14,10 @@
 
 #include "src/tint/sem/constant.h"
 
-#include <cmath>
-#include <utility>
-
-#include "src/tint/debug.h"
-#include "src/tint/program_builder.h"
-#include "src/tint/sem/type.h"
-
 namespace tint::sem {
 
-namespace {
-size_t CountElements(const Constant::Elements& elements) {
-    return std::visit([](auto&& vec) { return vec.size(); }, elements);
-}
-
-template <typename T>
-bool IsNegativeFloat(T value) {
-    (void)value;
-    if constexpr (IsFloatingPoint<T>) {
-        return std::signbit(value);
-    } else {
-        return false;
-    }
-}
-
-}  // namespace
-
-Constant::Constant() {}
-
-Constant::Constant(const sem::Type* ty, Elements els)
-    : type_(ty), elem_type_(CheckElemType(ty, CountElements(els))), elems_(std::move(els)) {}
-
-Constant::Constant(const sem::Type* ty, AInts vec) : Constant(ty, Elements{std::move(vec)}) {}
-
-Constant::Constant(const sem::Type* ty, AFloats vec) : Constant(ty, Elements{std::move(vec)}) {}
-
-Constant::Constant(const Constant&) = default;
+Constant::Constant() = default;
 
 Constant::~Constant() = default;
 
-Constant& Constant::operator=(const Constant& rhs) = default;
-
-bool Constant::AnyZero() const {
-    return WithElements([&](auto&& vec) {
-        using T = typename std::decay_t<decltype(vec)>::value_type;
-        for (auto el : vec) {
-            if (el == T(0) && !IsNegativeFloat(el.value)) {
-                return true;
-            }
-        }
-        return false;
-    });
-}
-
-bool Constant::AllZero(size_t start, size_t end) const {
-    return WithElements([&](auto&& vec) {
-        using T = typename std::decay_t<decltype(vec)>::value_type;
-        for (size_t i = start; i < end; i++) {
-            auto el = vec[i];
-            if (el != T(0) || IsNegativeFloat(el.value)) {
-                return false;
-            }
-        }
-        return true;
-    });
-}
-
-bool Constant::AllEqual(size_t start, size_t end) const {
-    return WithElements([&](auto&& vec) {
-        if (!vec.empty()) {
-            auto value = vec[start];
-            bool float_sign = IsNegativeFloat(vec[start].value);
-            for (size_t i = start + 1; i < end; i++) {
-                if (vec[i] != value || float_sign != IsNegativeFloat(vec[i].value)) {
-                    return false;
-                }
-            }
-        }
-        return true;
-    });
-}
-
-const Type* Constant::CheckElemType(const sem::Type* ty, size_t num_elements) {
-    diag::List diag;
-    uint32_t count = 0;
-    auto* el_ty = Type::DeepestElementOf(ty, &count);
-    if (!el_ty) {
-        TINT_ICE(Semantic, diag) << "Unsupported sem::Constant type: " << ty->TypeInfo().name;
-        return nullptr;
-    }
-    if (num_elements != count) {
-        TINT_ICE(Semantic, diag) << "sem::Constant() type <-> element mismatch. type: '"
-                                 << ty->TypeInfo().name << "' provided: " << num_elements
-                                 << " require: " << count;
-    }
-    TINT_ASSERT(Semantic, el_ty->is_abstract_or_scalar());
-    return el_ty;
-}
-
 }  // namespace tint::sem
diff --git a/src/tint/sem/constant.h b/src/tint/sem/constant.h
index cd73a79..b46127c 100644
--- a/src/tint/sem/constant.h
+++ b/src/tint/sem/constant.h
@@ -15,10 +15,7 @@
 #ifndef SRC_TINT_SEM_CONSTANT_H_
 #define SRC_TINT_SEM_CONSTANT_H_
 
-#include <ostream>
-#include <utility>
 #include <variant>
-#include <vector>
 
 #include "src/tint/number.h"
 
@@ -29,167 +26,53 @@
 
 namespace tint::sem {
 
-/// A Constant holds a compile-time evaluated expression value, expressed as a flattened list of
-/// element values. The expression type may be of an abstract-numeric, scalar, vector or matrix
-/// type. Constant holds the element values in either a vector of abstract-integer (AInt) or
-/// abstract-float (AFloat), depending on the element type.
+/// Constant is the interface to a compile-time evaluated expression value.
 class Constant {
   public:
-    /// AInts is a vector of AInt, used to hold elements of the WGSL types:
-    /// * abstract-integer
-    /// * i32
-    /// * u32
-    /// * bool (0 or 1)
-    using AInts = std::vector<AInt>;
-
-    /// AFloats is a vector of AFloat, used to hold elements of the WGSL types:
-    /// * abstract-float
-    /// * f32
-    /// * f16
-    using AFloats = std::vector<AFloat>;
-
-    /// Elements is either a vector of AInts or AFloats
-    using Elements = std::variant<AInts, AFloats>;
-
-    /// Helper that resolves to either AInt or AFloat based on the element type T.
-    template <typename T>
-    using ElementFor = std::conditional_t<IsFloatingPoint<UnwrapNumber<T>>, AFloat, AInt>;
-
-    /// Helper that resolves to either AInts or AFloats based on the element type T.
-    template <typename T>
-    using ElementVectorFor = std::conditional_t<IsFloatingPoint<UnwrapNumber<T>>, AFloats, AInts>;
-
-    /// Constructs an invalid Constant
+    /// Constructor
     Constant();
 
-    /// Constructs a Constant of the given type and element values
-    /// @param ty the Constant type
-    /// @param els the Constant element values
-    Constant(const sem::Type* ty, Elements els);
-
-    /// Constructs a Constant of the given type and element values
-    /// @param ty the Constant type
-    /// @param vec the Constant element values
-    Constant(const sem::Type* ty, AInts vec);
-
-    /// Constructs a Constant of the given type and element values
-    /// @param ty the Constant type
-    /// @param vec the Constant element values
-    Constant(const sem::Type* ty, AFloats vec);
-
-    /// Constructs a Constant of the given type and element values
-    /// @param ty the Constant type
-    /// @param els the Constant element values
-    template <typename T>
-    Constant(const sem::Type* ty, std::initializer_list<T> els);
-
-    /// Copy constructor
-    Constant(const Constant&);
-
     /// Destructor
-    ~Constant();
+    virtual ~Constant();
 
-    /// Copy assignment
-    /// @param other the Constant to copy
-    /// @returns this Constant
-    Constant& operator=(const Constant& other);
+    /// @returns the type of the constant
+    virtual const sem::Type* Type() const = 0;
 
-    /// @returns true if the Constant has been initialized
-    bool IsValid() const { return type_ != nullptr; }
+    /// @returns the value of this Constant, if this constant is of a scalar value or abstract
+    /// numeric, otherwsie std::monostate.
+    virtual std::variant<std::monostate, AInt, AFloat> Value() const = 0;
 
-    /// @return true if the Constant has been initialized
-    operator bool() const { return IsValid(); }
+    /// @returns the child constant element with the given index, or nullptr if the constant has no
+    /// children, or the index is out of bounds.
+    virtual const Constant* Index(size_t) const = 0;
 
-    /// @returns the type of the Constant
-    const sem::Type* Type() const { return type_; }
+    /// @returns true if child elements of this constant are positive-zero valued.
+    virtual bool AllZero() const = 0;
 
-    /// @returns the number of elements
-    size_t ElementCount() const {
-        return std::visit([](auto&& v) { return v.size(); }, elems_);
-    }
+    /// @returns true if any child elements of this constant are positive-zero valued.
+    virtual bool AnyZero() const = 0;
 
-    /// @returns the flattened element type of the Constant
-    const sem::Type* ElementType() const { return elem_type_; }
+    /// @returns true if all child elements of this constant have the same value and type.
+    virtual bool AllEqual() const = 0;
 
-    /// @returns the constant's flattened elements
-    const Elements& GetElements() const { return elems_; }
+    /// @returns a hash of the constant.
+    virtual size_t Hash() const = 0;
 
-    /// WithElements calls the function `f` with the vector of elements as either AFloats or AInts
-    /// @param f a function-like with the signature `R(auto&&)`.
-    /// @returns the result of calling `f`.
-    template <typename F>
-    auto WithElements(F&& f) const {
-        return std::visit(std::forward<F>(f), elems_);
-    }
-
-    /// WithElements calls the function `f` with the element vector as either AFloats or AInts
-    /// @param f a function-like with the signature `R(auto&&)`.
-    /// @returns the result of calling `f`.
-    template <typename F>
-    auto WithElements(F&& f) {
-        return std::visit(std::forward<F>(f), elems_);
-    }
-
-    /// @returns the elements as a vector of AInt
-    inline const AInts& IElements() const { return std::get<AInts>(elems_); }
-
-    /// @returns the elements as a vector of AFloat
-    inline const AFloats& FElements() const { return std::get<AFloats>(elems_); }
-
-    /// @returns true if any element is positive zero
-    bool AnyZero() const;
-
-    /// @returns true if all elements are positive zero
-    bool AllZero() const { return AllZero(0, ElementCount()); }
-
-    /// @returns true if all elements are the same value, with the same sign-bit.
-    bool AllEqual() const { return AllEqual(0, ElementCount()); }
-
-    /// @param start the first element index
-    /// @param end one past the last element index
-    /// @returns true if all elements between `[start, end)` are zero
-    bool AllZero(size_t start, size_t end) const;
-
-    /// @param start the first element index
-    /// @param end one past the last element index
-    /// @returns true if all elements between `[start, end)` are the same value
-    bool AllEqual(size_t start, size_t end) const;
-
-    /// @param index the index of the element
-    /// @return the element at `index`, which must be of type `T`.
+    /// @returns the value of the constant as the given scalar or abstract value.
     template <typename T>
-    T Element(size_t index) const;
-
-  private:
-    /// Checks that the provided type matches the number of expected elements.
-    /// @returns the element type of `ty`.
-    const sem::Type* CheckElemType(const sem::Type* ty, size_t num_elements);
-
-    const sem::Type* type_ = nullptr;
-    const sem::Type* elem_type_ = nullptr;
-    Elements elems_;
+    T As() const {
+        return std::visit(
+            [](auto v) {
+                if constexpr (std::is_same_v<decltype(v), std::monostate>) {
+                    return T(0);
+                } else {
+                    return static_cast<T>(v);
+                }
+            },
+            Value());
+    }
 };
 
-template <typename T>
-Constant::Constant(const sem::Type* ty, std::initializer_list<T> els)
-    : type_(ty), elem_type_(CheckElemType(type_, els.size())) {
-    ElementVectorFor<T> elements;
-    elements.reserve(els.size());
-    for (auto el : els) {
-        elements.emplace_back(ElementFor<T>(el));
-    }
-    elems_ = Elements{std::move(elements)};
-}
-
-template <typename T>
-T Constant::Element(size_t index) const {
-    if constexpr (std::is_same_v<ElementVectorFor<T>, AFloats>) {
-        return static_cast<T>(FElements()[index].value);
-    } else {
-        return static_cast<T>(IElements()[index].value);
-    }
-}
-
 }  // namespace tint::sem
 
 #endif  // SRC_TINT_SEM_CONSTANT_H_
diff --git a/src/tint/sem/constant_test.cc b/src/tint/sem/constant_test.cc
deleted file mode 100644
index e47e9a1..0000000
--- a/src/tint/sem/constant_test.cc
+++ /dev/null
@@ -1,540 +0,0 @@
-// 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/sem/constant.h"
-
-#include <gmock/gmock.h>
-
-#include "src/tint/sem/abstract_float.h"
-#include "src/tint/sem/abstract_int.h"
-#include "src/tint/sem/test_helper.h"
-
-using namespace tint::number_suffixes;  // NOLINT
-
-namespace tint::sem {
-namespace {
-
-struct ConstantTest : public TestHelper {
-    const sem::Array* Array(uint32_t n, const sem::Type* el_ty) {
-        return create<sem::Array>(el_ty,
-                                  /* count */ n,
-                                  /* align */ 16u,
-                                  /* size */ 4u * n,
-                                  /* stride */ 16u * n,
-                                  /* implicit_stride */ 16u * n);
-    }
-};
-
-TEST_F(ConstantTest, ConstructorInitializerList) {
-    {
-        auto i = AInt(AInt::kHighest);
-        Constant c(create<AbstractInt>(), {i});
-        c.WithElements([&](auto&& vec) { EXPECT_THAT(vec, testing::ElementsAre(i)); });
-    }
-    {
-        auto i = i32(i32::kHighest);
-        Constant c(create<I32>(), {i});
-        c.WithElements([&](auto&& vec) { EXPECT_THAT(vec, testing::ElementsAre(i)); });
-    }
-    {
-        auto i = u32(u32::kHighest);
-        Constant c(create<U32>(), {i});
-        c.WithElements([&](auto&& vec) { EXPECT_THAT(vec, testing::ElementsAre(i)); });
-    }
-    {
-        Constant c(create<Bool>(), {false});
-        c.WithElements([&](auto&& vec) { EXPECT_THAT(vec, testing::ElementsAre(0_a)); });
-    }
-    {
-        Constant c(create<Bool>(), {true});
-        c.WithElements([&](auto&& vec) { EXPECT_THAT(vec, testing::ElementsAre(1_a)); });
-    }
-    {
-        auto f = AFloat(AFloat::kHighest);
-        Constant c(create<AbstractFloat>(), {f});
-        c.WithElements([&](auto&& vec) { EXPECT_THAT(vec, testing::ElementsAre(f)); });
-    }
-    {
-        auto f = f32(f32::kHighest);
-        Constant c(create<F32>(), {f});
-        c.WithElements([&](auto&& vec) { EXPECT_THAT(vec, testing::ElementsAre(f)); });
-    }
-    {
-        auto f = f16(f16::kHighest);
-        Constant c(create<F16>(), {f});
-        c.WithElements([&](auto&& vec) { EXPECT_THAT(vec, testing::ElementsAre(f)); });
-    }
-}
-
-TEST_F(ConstantTest, Element_ai) {
-    auto* ty = create<AbstractInt>();
-    Constant c(ty, {1_a});
-    EXPECT_EQ(c.Element<AInt>(0), 1_a);
-    EXPECT_EQ(c.ElementCount(), 1u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), ty);
-}
-
-TEST_F(ConstantTest, Element_i32) {
-    auto* ty = create<I32>();
-    Constant c(ty, {1_a});
-    EXPECT_EQ(c.Element<i32>(0), 1_i);
-    EXPECT_EQ(c.ElementCount(), 1u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), ty);
-}
-
-TEST_F(ConstantTest, Element_u32) {
-    auto* ty = create<U32>();
-    Constant c(ty, {1_a});
-    EXPECT_EQ(c.Element<u32>(0), 1_u);
-    EXPECT_EQ(c.ElementCount(), 1u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), ty);
-}
-
-TEST_F(ConstantTest, Element_bool) {
-    auto* ty = create<Bool>();
-    Constant c(ty, {true});
-    EXPECT_EQ(c.Element<bool>(0), true);
-    EXPECT_EQ(c.ElementCount(), 1u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), ty);
-}
-
-TEST_F(ConstantTest, Element_af) {
-    auto* ty = create<AbstractFloat>();
-    Constant c(ty, {1.0_a});
-    EXPECT_EQ(c.Element<AFloat>(0), 1.0_a);
-    EXPECT_EQ(c.ElementCount(), 1u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), ty);
-}
-
-TEST_F(ConstantTest, Element_f32) {
-    auto* ty = create<F32>();
-    Constant c(ty, {1.0_a});
-    EXPECT_EQ(c.Element<f32>(0), 1.0_f);
-    EXPECT_EQ(c.ElementCount(), 1u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), ty);
-}
-
-TEST_F(ConstantTest, Element_f16) {
-    auto* ty = create<F16>();
-    Constant c(ty, {1.0_a});
-    EXPECT_EQ(c.Element<f16>(0), 1.0_h);
-    EXPECT_EQ(c.ElementCount(), 1u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), ty);
-}
-
-TEST_F(ConstantTest, Element_vec3_ai) {
-    auto* el_ty = create<AbstractInt>();
-    auto* ty = create<Vector>(el_ty, 3u);
-    Constant c(ty, {1_a, 2_a, 3_a});
-    EXPECT_EQ(c.Element<AInt>(0), 1_a);
-    EXPECT_EQ(c.Element<AInt>(1), 2_a);
-    EXPECT_EQ(c.Element<AInt>(2), 3_a);
-    EXPECT_EQ(c.ElementCount(), 3u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_vec3_i32) {
-    auto* el_ty = create<I32>();
-    auto* ty = create<Vector>(el_ty, 3u);
-    Constant c(ty, {1_a, 2_a, 3_a});
-    EXPECT_EQ(c.Element<i32>(0), 1_i);
-    EXPECT_EQ(c.Element<i32>(1), 2_i);
-    EXPECT_EQ(c.Element<i32>(2), 3_i);
-    EXPECT_EQ(c.ElementCount(), 3u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_vec3_u32) {
-    auto* el_ty = create<U32>();
-    auto* ty = create<Vector>(el_ty, 3u);
-    Constant c(ty, {1_a, 2_a, 3_a});
-    EXPECT_EQ(c.Element<u32>(0), 1_u);
-    EXPECT_EQ(c.Element<u32>(1), 2_u);
-    EXPECT_EQ(c.Element<u32>(2), 3_u);
-    EXPECT_EQ(c.ElementCount(), 3u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_vec3_bool) {
-    auto* el_ty = create<Bool>();
-    auto* ty = create<Vector>(el_ty, 2u);
-    Constant c(ty, {true, false});
-    EXPECT_EQ(c.Element<bool>(0), true);
-    EXPECT_EQ(c.Element<bool>(1), false);
-    EXPECT_EQ(c.ElementCount(), 2u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_vec3_af) {
-    auto* el_ty = create<AbstractFloat>();
-    auto* ty = create<Vector>(el_ty, 3u);
-    Constant c(ty, {1.0_a, 2.0_a, 3.0_a});
-    EXPECT_EQ(c.Element<AFloat>(0), 1.0_a);
-    EXPECT_EQ(c.Element<AFloat>(1), 2.0_a);
-    EXPECT_EQ(c.Element<AFloat>(2), 3.0_a);
-    EXPECT_EQ(c.ElementCount(), 3u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_vec3_f32) {
-    auto* el_ty = create<F32>();
-    auto* ty = create<Vector>(el_ty, 3u);
-    Constant c(ty, {1.0_a, 2.0_a, 3.0_a});
-    EXPECT_EQ(c.Element<f32>(0), 1.0_f);
-    EXPECT_EQ(c.Element<f32>(1), 2.0_f);
-    EXPECT_EQ(c.Element<f32>(2), 3.0_f);
-    EXPECT_EQ(c.ElementCount(), 3u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_vec3_f16) {
-    auto* el_ty = create<F16>();
-    auto* ty = create<Vector>(el_ty, 3u);
-    Constant c(ty, {1.0_a, 2.0_a, 3.0_a});
-    EXPECT_EQ(c.Element<f16>(0), 1.0_h);
-    EXPECT_EQ(c.Element<f16>(1), 2.0_h);
-    EXPECT_EQ(c.Element<f16>(2), 3.0_h);
-    EXPECT_EQ(c.ElementCount(), 3u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_mat2x3_af) {
-    auto* el_ty = create<AbstractFloat>();
-    auto* ty = create<Matrix>(create<Vector>(el_ty, 3u), 2u);
-    Constant c(ty, {1.0_a, 2.0_a, 3.0_a, 4.0_a, 5.0_a, 6.0_a});
-    EXPECT_EQ(c.Element<AFloat>(0), 1.0_a);
-    EXPECT_EQ(c.Element<AFloat>(1), 2.0_a);
-    EXPECT_EQ(c.Element<AFloat>(2), 3.0_a);
-    EXPECT_EQ(c.Element<AFloat>(3), 4.0_a);
-    EXPECT_EQ(c.Element<AFloat>(4), 5.0_a);
-    EXPECT_EQ(c.Element<AFloat>(5), 6.0_a);
-    EXPECT_EQ(c.ElementCount(), 6u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_mat2x3_f32) {
-    auto* el_ty = create<F32>();
-    auto* ty = create<Matrix>(create<Vector>(el_ty, 3u), 2u);
-    Constant c(ty, {1.0_a, 2.0_a, 3.0_a, 4.0_a, 5.0_a, 6.0_a});
-    EXPECT_EQ(c.Element<f32>(0), 1.0_f);
-    EXPECT_EQ(c.Element<f32>(1), 2.0_f);
-    EXPECT_EQ(c.Element<f32>(2), 3.0_f);
-    EXPECT_EQ(c.Element<f32>(3), 4.0_f);
-    EXPECT_EQ(c.Element<f32>(4), 5.0_f);
-    EXPECT_EQ(c.Element<f32>(5), 6.0_f);
-    EXPECT_EQ(c.ElementCount(), 6u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_mat2x3_f16) {
-    auto* el_ty = create<F16>();
-    auto* ty = create<Matrix>(create<Vector>(el_ty, 3u), 2u);
-    Constant c(ty, {1.0_a, 2.0_a, 3.0_a, 4.0_a, 5.0_a, 6.0_a});
-    EXPECT_EQ(c.Element<f16>(0), 1.0_h);
-    EXPECT_EQ(c.Element<f16>(1), 2.0_h);
-    EXPECT_EQ(c.Element<f16>(2), 3.0_h);
-    EXPECT_EQ(c.Element<f16>(3), 4.0_h);
-    EXPECT_EQ(c.Element<f16>(4), 5.0_h);
-    EXPECT_EQ(c.Element<f16>(5), 6.0_h);
-    EXPECT_EQ(c.ElementCount(), 6u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_arr_vec3_ai) {
-    auto* el_ty = create<AbstractInt>();
-    auto* ty = Array(2, create<Vector>(el_ty, 3u));
-    Constant c(ty, {1_a, 2_a, 3_a, 4_a, 5_a, 6_a});
-    EXPECT_EQ(c.Element<AInt>(0), 1_a);
-    EXPECT_EQ(c.Element<AInt>(1), 2_a);
-    EXPECT_EQ(c.Element<AInt>(2), 3_a);
-    EXPECT_EQ(c.Element<AInt>(3), 4_a);
-    EXPECT_EQ(c.Element<AInt>(4), 5_a);
-    EXPECT_EQ(c.Element<AInt>(5), 6_a);
-    EXPECT_EQ(c.ElementCount(), 6u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_arr_vec3_i32) {
-    auto* el_ty = create<I32>();
-    auto* ty = Array(2, create<Vector>(el_ty, 3u));
-    Constant c(ty, {1_a, 2_a, 3_a, 4_a, 5_a, 6_a});
-    EXPECT_EQ(c.Element<i32>(0), 1_i);
-    EXPECT_EQ(c.Element<i32>(1), 2_i);
-    EXPECT_EQ(c.Element<i32>(2), 3_i);
-    EXPECT_EQ(c.Element<i32>(3), 4_i);
-    EXPECT_EQ(c.Element<i32>(4), 5_i);
-    EXPECT_EQ(c.Element<i32>(5), 6_i);
-    EXPECT_EQ(c.ElementCount(), 6u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_arr_vec3_u32) {
-    auto* el_ty = create<U32>();
-    auto* ty = Array(2, create<Vector>(el_ty, 3u));
-    Constant c(ty, {1_a, 2_a, 3_a, 4_a, 5_a, 6_a});
-    EXPECT_EQ(c.Element<u32>(0), 1_u);
-    EXPECT_EQ(c.Element<u32>(1), 2_u);
-    EXPECT_EQ(c.Element<u32>(2), 3_u);
-    EXPECT_EQ(c.Element<u32>(3), 4_u);
-    EXPECT_EQ(c.Element<u32>(4), 5_u);
-    EXPECT_EQ(c.Element<u32>(5), 6_u);
-    EXPECT_EQ(c.ElementCount(), 6u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_arr_vec3_bool) {
-    auto* el_ty = create<Bool>();
-    auto* ty = Array(2, create<Vector>(el_ty, 2u));
-    Constant c(ty, {true, false, false, true});
-    EXPECT_EQ(c.Element<bool>(0), true);
-    EXPECT_EQ(c.Element<bool>(1), false);
-    EXPECT_EQ(c.Element<bool>(2), false);
-    EXPECT_EQ(c.Element<bool>(3), true);
-    EXPECT_EQ(c.ElementCount(), 4u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_arr_vec3_af) {
-    auto* el_ty = create<AbstractFloat>();
-    auto* ty = Array(2, create<Vector>(el_ty, 3u));
-    Constant c(ty, {1.0_a, 2.0_a, 3.0_a, 4.0_a, 5.0_a, 6.0_a});
-    EXPECT_EQ(c.Element<AFloat>(0), 1.0_a);
-    EXPECT_EQ(c.Element<AFloat>(1), 2.0_a);
-    EXPECT_EQ(c.Element<AFloat>(2), 3.0_a);
-    EXPECT_EQ(c.Element<AFloat>(3), 4.0_a);
-    EXPECT_EQ(c.Element<AFloat>(4), 5.0_a);
-    EXPECT_EQ(c.Element<AFloat>(5), 6.0_a);
-    EXPECT_EQ(c.ElementCount(), 6u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_arr_vec3_f32) {
-    auto* el_ty = create<F32>();
-    auto* ty = Array(2, create<Vector>(el_ty, 3u));
-    Constant c(ty, {1.0_a, 2.0_a, 3.0_a, 4.0_a, 5.0_a, 6.0_a});
-    EXPECT_EQ(c.Element<f32>(0), 1.0_f);
-    EXPECT_EQ(c.Element<f32>(1), 2.0_f);
-    EXPECT_EQ(c.Element<f32>(2), 3.0_f);
-    EXPECT_EQ(c.Element<f32>(3), 4.0_f);
-    EXPECT_EQ(c.Element<f32>(4), 5.0_f);
-    EXPECT_EQ(c.Element<f32>(5), 6.0_f);
-    EXPECT_EQ(c.ElementCount(), 6u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_arr_vec3_f16) {
-    auto* el_ty = create<F16>();
-    auto* ty = Array(2, create<Vector>(el_ty, 3u));
-    Constant c(ty, {1.0_a, 2.0_a, 3.0_a, 4.0_a, 5.0_a, 6.0_a});
-    EXPECT_EQ(c.Element<f16>(0), 1.0_h);
-    EXPECT_EQ(c.Element<f16>(1), 2.0_h);
-    EXPECT_EQ(c.Element<f16>(2), 3.0_h);
-    EXPECT_EQ(c.Element<f16>(3), 4.0_h);
-    EXPECT_EQ(c.Element<f16>(4), 5.0_h);
-    EXPECT_EQ(c.Element<f16>(5), 6.0_h);
-    EXPECT_EQ(c.ElementCount(), 6u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, Element_arr_arr_mat2x3_f32) {
-    auto* el_ty = create<F32>();
-    auto* ty = Array(2, Array(2, create<Matrix>(create<Vector>(el_ty, 3u), 2u)));
-    Constant c(ty, {
-                       1.0_a,  2.0_a,  3.0_a,  //
-                       4.0_a,  5.0_a,  6.0_a,  //
-
-                       7.0_a,  8.0_a,  9.0_a,   //
-                       10.0_a, 11.0_a, 12.0_a,  //
-
-                       13.0_a, 14.0_a, 15.0_a,  //
-                       16.0_a, 17.0_a, 18.0_a,  //
-
-                       19.0_a, 20.0_a, 21.0_a,  //
-                       22.0_a, 23.0_a, 24.0_a,  //
-                   });
-    for (size_t i = 0; i < 24; i++) {
-        EXPECT_EQ(c.Element<f32>(i), f32(i + 1));
-    }
-    EXPECT_EQ(c.ElementCount(), 24u);
-    EXPECT_TYPE(c.Type(), ty);
-    EXPECT_TYPE(c.ElementType(), el_ty);
-}
-
-TEST_F(ConstantTest, AnyZero) {
-    auto* vec3_ai = create<Vector>(create<AbstractInt>(), 3u);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 2_a, 3_a}).AnyZero(), false);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 2_a, 3_a}).AnyZero(), true);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 0_a, 3_a}).AnyZero(), true);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 2_a, 0_a}).AnyZero(), true);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 0_a, 0_a}).AnyZero(), true);
-
-    auto* vec3_af = create<Vector>(create<AbstractFloat>(), 3u);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 2._a, 3._a}).AnyZero(), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 2._a, 3._a}).AnyZero(), true);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 0._a, 3._a}).AnyZero(), true);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 2._a, 0._a}).AnyZero(), true);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 0._a, 0._a}).AnyZero(), true);
-
-    EXPECT_EQ(Constant(vec3_af, {1._a, -2._a, 3._a}).AnyZero(), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, -2._a, 3._a}).AnyZero(), true);
-    EXPECT_EQ(Constant(vec3_af, {1._a, -0._a, 3._a}).AnyZero(), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, -2._a, 0._a}).AnyZero(), true);
-    EXPECT_EQ(Constant(vec3_af, {0._a, -0._a, 0._a}).AnyZero(), true);
-    EXPECT_EQ(Constant(vec3_af, {-0._a, -0._a, -0._a}).AnyZero(), false);
-}
-
-TEST_F(ConstantTest, AllZero) {
-    auto* vec3_ai = create<Vector>(create<AbstractInt>(), 3u);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 2_a, 3_a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 2_a, 3_a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 0_a, 3_a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 2_a, 0_a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 0_a, 0_a}).AllZero(), true);
-
-    auto* vec3_af = create<Vector>(create<AbstractFloat>(), 3u);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 2._a, 3._a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 2._a, 3._a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 0._a, 3._a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 2._a, 0._a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 0._a, 0._a}).AllZero(), true);
-
-    EXPECT_EQ(Constant(vec3_af, {1._a, -2._a, 3._a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, -2._a, 3._a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, -0._a, 3._a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, -2._a, 0._a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, -0._a, 0._a}).AllZero(), false);
-    EXPECT_EQ(Constant(vec3_af, {-0._a, -0._a, -0._a}).AllZero(), false);
-}
-
-TEST_F(ConstantTest, AllEqual) {
-    auto* vec3_ai = create<Vector>(create<AbstractInt>(), 3u);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 2_a, 3_a}).AllEqual(), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 1_a, 3_a}).AllEqual(), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 3_a, 3_a}).AllEqual(), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 1_a, 1_a}).AllEqual(), true);
-    EXPECT_EQ(Constant(vec3_ai, {2_a, 2_a, 2_a}).AllEqual(), true);
-    EXPECT_EQ(Constant(vec3_ai, {3_a, 3_a, 3_a}).AllEqual(), true);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 0_a, 0_a}).AllEqual(), true);
-
-    auto* vec3_af = create<Vector>(create<AbstractFloat>(), 3u);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 2._a, 3._a}).AllEqual(), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 1._a, 3._a}).AllEqual(), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 3._a, 3._a}).AllEqual(), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 1._a, 1._a}).AllEqual(), true);
-    EXPECT_EQ(Constant(vec3_af, {2._a, 2._a, 2._a}).AllEqual(), true);
-    EXPECT_EQ(Constant(vec3_af, {3._a, 3._a, 3._a}).AllEqual(), true);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 0._a, 0._a}).AllEqual(), true);
-    EXPECT_EQ(Constant(vec3_af, {0._a, -0._a, 0._a}).AllEqual(), false);
-}
-
-TEST_F(ConstantTest, AllZeroRange) {
-    auto* vec3_ai = create<Vector>(create<AbstractInt>(), 3u);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 2_a, 3_a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 2_a, 3_a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 2_a, 3_a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 2_a, 0_a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 0_a, 3_a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 2_a, 0_a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 0_a, 0_a}).AllZero(1, 3), true);
-
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 2_a, 3_a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 2_a, 3_a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 2_a, 3_a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 2_a, 0_a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 0_a, 3_a}).AllZero(0, 2), true);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 2_a, 0_a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 0_a, 0_a}).AllZero(0, 2), false);
-
-    auto* vec3_af = create<Vector>(create<AbstractFloat>(), 3u);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 2._a, 3._a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 2._a, 3._a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 2._a, 3._a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 2._a, 0._a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 0._a, 3._a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 2._a, 0._a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 0._a, 0._a}).AllZero(1, 3), true);
-    EXPECT_EQ(Constant(vec3_af, {1._a, -0._a, 0._a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 0._a, -0._a}).AllZero(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, -0._a, -0._a}).AllZero(1, 3), false);
-
-    EXPECT_EQ(Constant(vec3_af, {1._a, 2._a, 3._a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 2._a, 3._a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 2._a, 3._a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 2._a, 0._a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 0._a, 3._a}).AllZero(0, 2), true);
-    EXPECT_EQ(Constant(vec3_af, {-0._a, 0._a, 1._a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, -0._a, 1._a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_af, {-0._a, -0._a, 1._a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 2._a, 0._a}).AllZero(0, 2), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 0._a, 0._a}).AllZero(0, 2), false);
-}
-
-TEST_F(ConstantTest, AllEqualRange) {
-    auto* vec3_ai = create<Vector>(create<AbstractInt>(), 3u);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 2_a, 3_a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 1_a, 3_a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 3_a, 3_a}).AllEqual(1, 3), true);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 1_a, 1_a}).AllEqual(1, 3), true);
-    EXPECT_EQ(Constant(vec3_ai, {2_a, 2_a, 2_a}).AllEqual(1, 3), true);
-    EXPECT_EQ(Constant(vec3_ai, {2_a, 2_a, 3_a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_ai, {1_a, 0_a, 0_a}).AllEqual(1, 3), true);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 1_a, 0_a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 0_a, 1_a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_ai, {0_a, 0_a, 0_a}).AllEqual(1, 3), true);
-
-    auto* vec3_af = create<Vector>(create<AbstractFloat>(), 3u);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 2._a, 3._a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 1._a, 3._a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 3._a, 3._a}).AllEqual(1, 3), true);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 1._a, 1._a}).AllEqual(1, 3), true);
-    EXPECT_EQ(Constant(vec3_af, {2._a, 2._a, 2._a}).AllEqual(1, 3), true);
-    EXPECT_EQ(Constant(vec3_af, {2._a, 2._a, 3._a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {1._a, 0._a, 0._a}).AllEqual(1, 3), true);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 1._a, 0._a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 0._a, 1._a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, 0._a, 0._a}).AllEqual(1, 3), true);
-    EXPECT_EQ(Constant(vec3_af, {1._a, -0._a, 0._a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, -1._a, 0._a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, -0._a, 1._a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, -0._a, 0._a}).AllEqual(1, 3), false);
-    EXPECT_EQ(Constant(vec3_af, {0._a, -0._a, -0._a}).AllEqual(1, 3), true);
-    EXPECT_EQ(Constant(vec3_af, {-0._a, -0._a, -0._a}).AllEqual(1, 3), true);
-}
-
-}  // namespace
-}  // namespace tint::sem
diff --git a/src/tint/sem/expression.cc b/src/tint/sem/expression.cc
index 57ec68b..4415db5 100644
--- a/src/tint/sem/expression.cc
+++ b/src/tint/sem/expression.cc
@@ -25,7 +25,7 @@
 Expression::Expression(const ast::Expression* declaration,
                        const sem::Type* type,
                        const Statement* statement,
-                       Constant constant,
+                       const Constant* constant,
                        bool has_side_effects,
                        const Variable* source_var /* = nullptr */)
     : declaration_(declaration),
diff --git a/src/tint/sem/expression.h b/src/tint/sem/expression.h
index a783851..05e5dac 100644
--- a/src/tint/sem/expression.h
+++ b/src/tint/sem/expression.h
@@ -41,7 +41,7 @@
     Expression(const ast::Expression* declaration,
                const sem::Type* type,
                const Statement* statement,
-               Constant constant,
+               const Constant* constant,
                bool has_side_effects,
                const Variable* source_var = nullptr);
 
@@ -58,7 +58,7 @@
     const Statement* Stmt() const { return statement_; }
 
     /// @return the constant value of this expression
-    const Constant& ConstantValue() const { return constant_; }
+    const Constant* ConstantValue() const { return constant_; }
 
     /// Returns the variable or parameter that this expression derives from.
     /// For reference and pointer expressions, this will either be the originating
@@ -88,7 +88,7 @@
   private:
     const sem::Type* const type_;
     const Statement* const statement_;
-    const Constant constant_;
+    const Constant* const constant_;
     sem::Behaviors behaviors_{sem::Behavior::kNext};
     const bool has_side_effects_;
 };
diff --git a/src/tint/sem/expression_test.cc b/src/tint/sem/expression_test.cc
index fc1adeb..cc4bf0e 100644
--- a/src/tint/sem/expression_test.cc
+++ b/src/tint/sem/expression_test.cc
@@ -23,13 +23,30 @@
 namespace tint::sem {
 namespace {
 
+class MockConstant : public sem::Constant {
+  public:
+    explicit MockConstant(const sem::Type* ty) : type(ty) {}
+    ~MockConstant() override {}
+    const sem::Type* Type() const override { return type; }
+    std::variant<std::monostate, AInt, AFloat> Value() const override { return {}; }
+    const Constant* Index(size_t) const override { return {}; }
+    bool AllZero() const override { return {}; }
+    bool AnyZero() const override { return {}; }
+    bool AllEqual() const override { return {}; }
+    size_t Hash() const override { return 0; }
+
+  private:
+    const sem::Type* type;
+};
+
 using ExpressionTest = TestHelper;
 
 TEST_F(ExpressionTest, UnwrapMaterialize) {
+    MockConstant c(create<I32>());
     auto* a = create<Expression>(/* declaration */ nullptr, create<I32>(), /* statement */ nullptr,
-                                 Constant{},
+                                 /* constant_value */ nullptr,
                                  /* has_side_effects */ false, /* source_var */ nullptr);
-    auto* b = create<Materialize>(a, /* statement */ nullptr, Constant{create<I32>(), {1_a}});
+    auto* b = create<Materialize>(a, /* statement */ nullptr, &c);
 
     EXPECT_EQ(a, a->UnwrapMaterialize());
     EXPECT_EQ(a, b->UnwrapMaterialize());
diff --git a/src/tint/sem/index_accessor_expression.cc b/src/tint/sem/index_accessor_expression.cc
index 06f0990..cd74201 100644
--- a/src/tint/sem/index_accessor_expression.cc
+++ b/src/tint/sem/index_accessor_expression.cc
@@ -27,7 +27,7 @@
                                                  const Expression* object,
                                                  const Expression* index,
                                                  const Statement* statement,
-                                                 Constant constant,
+                                                 const Constant* constant,
                                                  bool has_side_effects,
                                                  const Variable* source_var /* = nullptr */)
     : Base(declaration, type, statement, constant, has_side_effects, source_var),
diff --git a/src/tint/sem/index_accessor_expression.h b/src/tint/sem/index_accessor_expression.h
index c77f55d..233b0fa 100644
--- a/src/tint/sem/index_accessor_expression.h
+++ b/src/tint/sem/index_accessor_expression.h
@@ -43,7 +43,7 @@
                             const Expression* object,
                             const Expression* index,
                             const Statement* statement,
-                            Constant constant,
+                            const Constant* constant,
                             bool has_side_effects,
                             const Variable* source_var = nullptr);
 
diff --git a/src/tint/sem/materialize.cc b/src/tint/sem/materialize.cc
index 76dd9d4..739b0ce 100644
--- a/src/tint/sem/materialize.cc
+++ b/src/tint/sem/materialize.cc
@@ -17,19 +17,16 @@
 TINT_INSTANTIATE_TYPEINFO(tint::sem::Materialize);
 
 namespace tint::sem {
-
-Materialize::Materialize(const Expression* expr, const Statement* statement, Constant constant)
+Materialize::Materialize(const Expression* expr,
+                         const Statement* statement,
+                         const Constant* constant)
     : Base(/* declaration */ expr->Declaration(),
-           /* type */ constant.Type(),
+           /* type */ constant->Type(),
            /* statement */ statement,
            /* constant */ constant,
            /* has_side_effects */ false,
            /* source_var */ expr->SourceVariable()),
-      expr_(expr) {
-    // Materialize nodes only wrap compile-time expressions, and so the Materialize expression must
-    // have a constant value.
-    TINT_ASSERT(Semantic, constant.IsValid());
-}
+      expr_(expr) {}
 
 Materialize::~Materialize() = default;
 
diff --git a/src/tint/sem/materialize.h b/src/tint/sem/materialize.h
index a7c0e3a..0cbac29 100644
--- a/src/tint/sem/materialize.h
+++ b/src/tint/sem/materialize.h
@@ -31,7 +31,7 @@
     /// @param expr the inner expression, being materialized
     /// @param statement the statement that owns this expression
     /// @param constant the constant value of this expression
-    Materialize(const Expression* expr, const Statement* statement, Constant constant);
+    Materialize(const Expression* expr, const Statement* statement, const Constant* constant);
 
     /// Destructor
     ~Materialize() override;
diff --git a/src/tint/sem/member_accessor_expression.cc b/src/tint/sem/member_accessor_expression.cc
index bd706a5..f9929c7 100644
--- a/src/tint/sem/member_accessor_expression.cc
+++ b/src/tint/sem/member_accessor_expression.cc
@@ -29,8 +29,7 @@
                                                    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, statement, nullptr, has_side_effects, source_var), object_(object) {}
 
 MemberAccessorExpression::~MemberAccessorExpression() = default;
 
diff --git a/src/tint/sem/variable.cc b/src/tint/sem/variable.cc
index 5439600..3849807 100644
--- a/src/tint/sem/variable.cc
+++ b/src/tint/sem/variable.cc
@@ -33,7 +33,7 @@
                    const sem::Type* type,
                    ast::StorageClass storage_class,
                    ast::Access access,
-                   Constant constant_value)
+                   const Constant* constant_value)
     : declaration_(declaration),
       type_(type),
       storage_class_(storage_class),
@@ -47,9 +47,8 @@
                              ast::StorageClass storage_class,
                              ast::Access access,
                              const sem::Statement* statement,
-                             Constant constant_value)
-    : Base(declaration, type, storage_class, access, std::move(constant_value)),
-      statement_(statement) {}
+                             const Constant* constant_value)
+    : Base(declaration, type, storage_class, access, constant_value), statement_(statement) {}
 
 LocalVariable::~LocalVariable() = default;
 
@@ -57,9 +56,9 @@
                                const sem::Type* type,
                                ast::StorageClass storage_class,
                                ast::Access access,
-                               Constant constant_value,
+                               const Constant* constant_value,
                                sem::BindingPoint binding_point)
-    : Base(declaration, type, storage_class, access, std::move(constant_value)),
+    : Base(declaration, type, storage_class, access, constant_value),
       binding_point_(binding_point) {}
 
 GlobalVariable::~GlobalVariable() = default;
@@ -70,7 +69,7 @@
                      ast::StorageClass storage_class,
                      ast::Access access,
                      const ParameterUsage usage /* = ParameterUsage::kNone */)
-    : Base(declaration, type, storage_class, access, Constant{}), index_(index), usage_(usage) {}
+    : Base(declaration, type, storage_class, access, nullptr), index_(index), usage_(usage) {}
 
 Parameter::~Parameter() = default;
 
diff --git a/src/tint/sem/variable.h b/src/tint/sem/variable.h
index 3b5bd0c..28e8f97 100644
--- a/src/tint/sem/variable.h
+++ b/src/tint/sem/variable.h
@@ -52,7 +52,7 @@
              const sem::Type* type,
              ast::StorageClass storage_class,
              ast::Access access,
-             Constant constant_value);
+             const Constant* constant_value);
 
     /// Destructor
     ~Variable() override;
@@ -70,7 +70,7 @@
     ast::Access Access() const { return access_; }
 
     /// @return the constant value of this expression
-    const Constant& ConstantValue() const { return constant_value_; }
+    const Constant* ConstantValue() const { return constant_value_; }
 
     /// @returns the variable constructor expression, or nullptr if the variable
     /// does not have one.
@@ -91,7 +91,7 @@
     const sem::Type* const type_;
     const ast::StorageClass storage_class_;
     const ast::Access access_;
-    const Constant constant_value_;
+    const Constant* constant_value_;
     const Expression* constructor_ = nullptr;
     std::vector<const VariableUser*> users_;
 };
@@ -111,7 +111,7 @@
                   ast::StorageClass storage_class,
                   ast::Access access,
                   const sem::Statement* statement,
-                  Constant constant_value);
+                  const Constant* constant_value);
 
     /// Destructor
     ~LocalVariable() override;
@@ -145,7 +145,7 @@
                    const sem::Type* type,
                    ast::StorageClass storage_class,
                    ast::Access access,
-                   Constant constant_value,
+                   const Constant* constant_value,
                    sem::BindingPoint binding_point = {});
 
     /// Destructor
diff --git a/src/tint/transform/localize_struct_array_assignment.cc b/src/tint/transform/localize_struct_array_assignment.cc
index d6cdded..7c3d695 100644
--- a/src/tint/transform/localize_struct_array_assignment.cc
+++ b/src/tint/transform/localize_struct_array_assignment.cc
@@ -46,7 +46,7 @@
             expr, b.Diagnostics(), [&](const ast::IndexAccessorExpression* ia) {
                 // Indexing using a runtime value?
                 auto* idx_sem = ctx.src->Sem().Get(ia->index);
-                if (!idx_sem->ConstantValue().IsValid()) {
+                if (!idx_sem->ConstantValue()) {
                     // Indexing a member access expr?
                     if (auto* ma = ia->object->As<ast::MemberAccessorExpression>()) {
                         // That accesses an array?
diff --git a/src/tint/transform/promote_side_effects_to_decl.cc b/src/tint/transform/promote_side_effects_to_decl.cc
index d527a4c..7551c23 100644
--- a/src/tint/transform/promote_side_effects_to_decl.cc
+++ b/src/tint/transform/promote_side_effects_to_decl.cc
@@ -275,7 +275,7 @@
                 if (auto* sem_e = sem.Get(e)) {
                     if (auto* var_user = sem_e->As<sem::VariableUser>()) {
                         // Don't hoist constants.
-                        if (var_user->ConstantValue().IsValid()) {
+                        if (var_user->ConstantValue()) {
                             return false;
                         }
                         // Don't hoist read-only variables as they cannot receive
diff --git a/src/tint/transform/robustness.cc b/src/tint/transform/robustness.cc
index 2e1cc40..beb1108 100644
--- a/src/tint/transform/robustness.cc
+++ b/src/tint/transform/robustness.cc
@@ -120,17 +120,18 @@
             return nullptr;
         }
 
-        if (auto idx_constant = idx_sem->ConstantValue()) {
+        if (auto* idx_constant = idx_sem->ConstantValue()) {
             // Constant value index
-            if (idx_constant.Type()->Is<sem::I32>()) {
-                idx.i32 = static_cast<int32_t>(idx_constant.Element<AInt>(0).value);
+            auto val = std::get<AInt>(idx_constant->Value());
+            if (idx_constant->Type()->Is<sem::I32>()) {
+                idx.i32 = static_cast<int32_t>(val);
                 idx.is_signed = true;
-            } else if (idx_constant.Type()->Is<sem::U32>()) {
-                idx.u32 = static_cast<uint32_t>(idx_constant.Element<AInt>(0).value);
+            } else if (idx_constant->Type()->Is<sem::U32>()) {
+                idx.u32 = static_cast<uint32_t>(val);
                 idx.is_signed = false;
             } else {
                 TINT_ICE(Transform, b.Diagnostics()) << "unsupported constant value for accessor "
-                                                     << idx_constant.Type()->TypeInfo().name;
+                                                     << idx_constant->Type()->TypeInfo().name;
                 return nullptr;
             }
         } else {
diff --git a/src/tint/transform/zero_init_workgroup_memory.cc b/src/tint/transform/zero_init_workgroup_memory.cc
index f56dc61..96395f1 100644
--- a/src/tint/transform/zero_init_workgroup_memory.cc
+++ b/src/tint/transform/zero_init_workgroup_memory.cc
@@ -358,8 +358,8 @@
                 continue;
             }
             auto* sem = ctx.src->Sem().Get(expr);
-            if (auto c = sem->ConstantValue()) {
-                workgroup_size_const *= c.Element<AInt>(0).value;
+            if (auto* c = sem->ConstantValue()) {
+                workgroup_size_const *= c->As<AInt>();
                 continue;
             }
             // Constant value could not be found. Build expression instead.
diff --git a/src/tint/writer/append_vector.cc b/src/tint/writer/append_vector.cc
index 7755ae3..c5f184d 100644
--- a/src/tint/writer/append_vector.cc
+++ b/src/tint/writer/append_vector.cc
@@ -59,7 +59,7 @@
             << "unsupported vector element type: " << ty->TypeInfo().name;
         return nullptr;
     }
-    auto* sem = b.create<sem::Expression>(expr, ty, stmt, sem::Constant{},
+    auto* sem = b.create<sem::Expression>(expr, ty, stmt, /* constant_value */ nullptr,
                                           /* has_side_effects */ false);
     b.Sem().Add(expr, sem);
     return sem;
@@ -139,7 +139,7 @@
                                       ast::StorageClass::kNone, ast::Access::kUndefined));
         auto* scalar_cast_sem = b->create<sem::Call>(
             scalar_cast_ast, scalar_cast_target, std::vector<const sem::Expression*>{scalar_sem},
-            statement, sem::Constant{}, /* has_side_effects */ false);
+            statement, /* constant_value */ nullptr, /* has_side_effects */ false);
         b->Sem().Add(scalar_cast_ast, scalar_cast_sem);
         packed.emplace_back(scalar_cast_sem);
     } else {
@@ -158,7 +158,7 @@
                                                  ast::Access::kUndefined);
             }));
     auto* constructor_sem = b->create<sem::Call>(constructor_ast, constructor_target, packed,
-                                                 statement, sem::Constant{},
+                                                 statement, /* constant_value */ nullptr,
                                                  /* has_side_effects */ false);
     b->Sem().Add(constructor_ast, constructor_sem);
     return constructor_sem;
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index 022db4b..fdf4557 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -1332,7 +1332,7 @@
 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, sem::Constant{},
+    auto* sem_zero = builder_.create<sem::Expression>(zero, f32, stmt, /* constant_value */ nullptr,
                                                       /* has_side_effects */ false);
     builder_.Sem().Add(zero, sem_zero);
     return zero;
@@ -1771,7 +1771,7 @@
 
 bool GeneratorImpl::EmitExpression(std::ostream& out, const ast::Expression* expr) {
     if (auto* sem = builder_.Sem().Get(expr)) {
-        if (auto constant = sem->ConstantValue()) {
+        if (auto* constant = sem->ConstantValue()) {
             return EmitConstant(out, constant);
         }
     }
@@ -2214,31 +2214,23 @@
     return true;
 }
 
-bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant& constant) {
-    return EmitConstantRange(out, constant, constant.Type(), 0, constant.ElementCount());
-}
-
-bool GeneratorImpl::EmitConstantRange(std::ostream& out,
-                                      const sem::Constant& constant,
-                                      const sem::Type* range_ty,
-                                      size_t start,
-                                      size_t end) {
+bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constant) {
     return Switch(
-        range_ty,  //
+        constant->Type(),  //
         [&](const sem::Bool*) {
-            out << (constant.Element<AInt>(start) ? "true" : "false");
+            out << (constant->As<AInt>() ? "true" : "false");
             return true;
         },
         [&](const sem::F32*) {
-            PrintF32(out, static_cast<float>(constant.Element<AFloat>(start)));
+            PrintF32(out, constant->As<float>());
             return true;
         },
         [&](const sem::I32*) {
-            out << constant.Element<AInt>(start).value;
+            out << constant->As<AInt>();
             return true;
         },
         [&](const sem::U32*) {
-            out << constant.Element<AInt>(start).value << "u";
+            out << constant->As<AInt>() << "u";
             return true;
         },
         [&](const sem::Vector* v) {
@@ -2248,15 +2240,15 @@
 
             ScopedParen sp(out);
 
-            if (constant.AllEqual(start, end)) {
-                return EmitConstantRange(out, constant, v->type(), start, start + 1);
+            if (constant->AllEqual()) {
+                return EmitConstant(out, constant->Index(0));
             }
 
-            for (size_t i = start; i < end; i++) {
-                if (i > start) {
+            for (size_t i = 0; i < v->Width(); i++) {
+                if (i > 0) {
                     out << ", ";
                 }
-                if (!EmitConstantRange(out, constant, v->type(), i, i + 1u)) {
+                if (!EmitConstant(out, constant->Index(i))) {
                     return false;
                 }
             }
@@ -2273,9 +2265,7 @@
                 if (column_idx > 0) {
                     out << ", ";
                 }
-                size_t col_start = m->rows() * column_idx;
-                size_t col_end = col_start + m->rows();
-                if (!EmitConstantRange(out, constant, m->ColumnType(), col_start, col_end)) {
+                if (!EmitConstant(out, constant->Index(column_idx))) {
                     return false;
                 }
             }
@@ -2288,15 +2278,11 @@
 
             ScopedParen sp(out);
 
-            auto* el_ty = a->ElemType();
-
-            uint32_t step = 0;
-            sem::Type::DeepestElementOf(el_ty, &step);
-            for (size_t i = start; i < end; i += step) {
-                if (i > start) {
+            for (size_t i = 0; i < a->Count(); i++) {
+                if (i > 0) {
                     out << ", ";
                 }
-                if (!EmitConstantRange(out, constant, el_ty, i, i + step)) {
+                if (!EmitConstant(out, constant->Index(i))) {
                     return false;
                 }
             }
@@ -2306,7 +2292,7 @@
         [&](Default) {
             diagnostics_.add_error(
                 diag::System::Writer,
-                "unhandled constant type: " + builder_.FriendlyName(constant.Type()));
+                "unhandled constant type: " + builder_.FriendlyName(constant->Type()));
             return false;
         });
 }
diff --git a/src/tint/writer/glsl/generator_impl.h b/src/tint/writer/glsl/generator_impl.h
index f925bba..15b7ee1 100644
--- a/src/tint/writer/glsl/generator_impl.h
+++ b/src/tint/writer/glsl/generator_impl.h
@@ -346,19 +346,7 @@
     /// @param out the output stream
     /// @param constant the constant value to emit
     /// @returns true if the constant value was successfully emitted
-    bool EmitConstant(std::ostream& out, const sem::Constant& constant);
-    /// Handles emitting a sub-range of a constant value
-    /// @param out the output stream
-    /// @param constant the constant value to emit
-    /// @param range_ty the sub-range type
-    /// @param start the element index for the first element
-    /// @param end the element index for one past the last element
-    /// @returns true if the constant value was successfully emitted
-    bool EmitConstantRange(std::ostream& out,
-                           const sem::Constant& constant,
-                           const sem::Type* range_ty,
-                           size_t start,
-                           size_t end);
+    bool EmitConstant(std::ostream& out, const sem::Constant* constant);
     /// Handles a literal
     /// @param out the output stream
     /// @param lit the literal to emit
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index 3e4352b..b307c8d 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -615,8 +615,7 @@
             if (auto* mat = TypeOf(lhs_sub_access->object)->UnwrapRef()->As<sem::Matrix>()) {
                 auto* rhs_col_idx_sem = builder_.Sem().Get(lhs_access->index);
                 auto* rhs_row_idx_sem = builder_.Sem().Get(lhs_sub_access->index);
-                if (!rhs_col_idx_sem->ConstantValue().IsValid() ||
-                    !rhs_row_idx_sem->ConstantValue().IsValid()) {
+                if (!rhs_col_idx_sem->ConstantValue() || !rhs_row_idx_sem->ConstantValue()) {
                     return EmitDynamicMatrixScalarAssignment(stmt, mat);
                 }
             }
@@ -626,7 +625,7 @@
         const auto* lhs_access_type = TypeOf(lhs_access->object)->UnwrapRef();
         if (auto* mat = lhs_access_type->As<sem::Matrix>()) {
             auto* lhs_index_sem = builder_.Sem().Get(lhs_access->index);
-            if (!lhs_index_sem->ConstantValue().IsValid()) {
+            if (!lhs_index_sem->ConstantValue()) {
                 return EmitDynamicMatrixVectorAssignment(stmt, mat);
             }
         }
@@ -634,7 +633,7 @@
         // indices
         if (auto* vec = lhs_access_type->As<sem::Vector>()) {
             auto* rhs_sem = builder_.Sem().Get(lhs_access->index);
-            if (!rhs_sem->ConstantValue().IsValid()) {
+            if (!rhs_sem->ConstantValue()) {
                 return EmitDynamicVectorAssignment(stmt, vec);
             }
         }
@@ -654,28 +653,30 @@
 
 bool GeneratorImpl::EmitExpressionOrOneIfZero(std::ostream& out, const ast::Expression* expr) {
     // For constants, replace literal 0 with 1.
-    if (const auto& val = builder_.Sem().Get(expr)->ConstantValue()) {
-        if (!val.AnyZero()) {
+    if (const auto* val = builder_.Sem().Get(expr)->ConstantValue()) {
+        if (!val->AnyZero()) {
             return EmitExpression(out, expr);
         }
 
-        if (val.Type()->IsAnyOf<sem::I32, sem::U32>()) {
-            return EmitValue(out, val.Type(), 1);
+        auto* ty = val->Type();
+
+        if (ty->IsAnyOf<sem::I32, sem::U32>()) {
+            return EmitValue(out, ty, 1);
         }
 
-        if (auto* vec = val.Type()->As<sem::Vector>()) {
+        if (auto* vec = ty->As<sem::Vector>()) {
             auto* elem_ty = vec->type();
 
-            if (!EmitType(out, val.Type(), ast::StorageClass::kNone, ast::Access::kUndefined, "")) {
+            if (!EmitType(out, ty, ast::StorageClass::kNone, ast::Access::kUndefined, "")) {
                 return false;
             }
 
             out << "(";
-            for (size_t i = 0; i < val.ElementCount(); ++i) {
+            for (size_t i = 0; i < vec->Width(); ++i) {
                 if (i != 0) {
                     out << ", ";
                 }
-                auto s = val.Element<AInt>(i).value;
+                auto s = val->Index(i)->As<AInt>();
                 if (!EmitValue(out, elem_ty, (s == 0) ? 1 : static_cast<int>(s))) {
                     return false;
                 }
@@ -1181,9 +1182,9 @@
     // If true, use scalar_offset_value, otherwise use scalar_offset_expr
     bool scalar_offset_constant = false;
 
-    if (auto val = offset_arg->ConstantValue()) {
-        TINT_ASSERT(Writer, val.Type()->Is<sem::U32>());
-        scalar_offset_value = static_cast<uint32_t>(val.Element<AInt>(0).value);
+    if (auto* val = offset_arg->ConstantValue()) {
+        TINT_ASSERT(Writer, val->Type()->Is<sem::U32>());
+        scalar_offset_value = static_cast<uint32_t>(std::get<AInt>(val->Value()));
         scalar_offset_value /= 4;  // bytes -> scalar index
         scalar_offset_constant = true;
     }
@@ -2337,7 +2338,7 @@
         case sem::BuiltinType::kTextureGather:
             out << ".Gather";
             if (builtin->Parameters()[0]->Usage() == sem::ParameterUsage::kComponent) {
-                switch (call->Arguments()[0]->ConstantValue().Element<AInt>(0).value) {
+                switch (call->Arguments()[0]->ConstantValue()->As<AInt>()) {
                     case 0:
                         out << "Red";
                         break;
@@ -2384,8 +2385,9 @@
         auto* i32 = builder_.create<sem::I32>();
         auto* zero = builder_.Expr(0_i);
         auto* stmt = builder_.Sem().Get(vector)->Stmt();
-        builder_.Sem().Add(zero, builder_.create<sem::Expression>(zero, i32, stmt, sem::Constant{},
-                                                                  /* has_side_effects */ false));
+        builder_.Sem().Add(
+            zero, builder_.create<sem::Expression>(zero, i32, stmt, /* constant_value */ nullptr,
+                                                   /* has_side_effects */ false));
         auto* packed = AppendVector(&builder_, vector, zero);
         return EmitExpression(out, packed->Declaration());
     };
@@ -2614,7 +2616,7 @@
 
 bool GeneratorImpl::EmitExpression(std::ostream& out, const ast::Expression* expr) {
     if (auto* sem = builder_.Sem().Get(expr)) {
-        if (auto constant = sem->ConstantValue()) {
+        if (auto* constant = sem->ConstantValue()) {
             return EmitConstant(out, constant);
         }
     }
@@ -3109,43 +3111,35 @@
     return true;
 }
 
-bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant& constant) {
-    return EmitConstantRange(out, constant, constant.Type(), 0, constant.ElementCount());
-}
-
-bool GeneratorImpl::EmitConstantRange(std::ostream& out,
-                                      const sem::Constant& constant,
-                                      const sem::Type* range_ty,
-                                      size_t start,
-                                      size_t end) {
+bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constant) {
     return Switch(
-        range_ty,  //
+        constant->Type(),  //
         [&](const sem::Bool*) {
-            out << (constant.Element<AInt>(start) ? "true" : "false");
+            out << (constant->As<AInt>() ? "true" : "false");
             return true;
         },
         [&](const sem::F32*) {
-            PrintF32(out, static_cast<float>(constant.Element<AFloat>(start)));
+            PrintF32(out, constant->As<float>());
             return true;
         },
         [&](const sem::I32*) {
-            out << constant.Element<AInt>(start).value;
+            out << constant->As<AInt>();
             return true;
         },
         [&](const sem::U32*) {
-            out << constant.Element<AInt>(start).value << "u";
+            out << constant->As<AInt>() << "u";
             return true;
         },
         [&](const sem::Vector* v) {
-            if (constant.AllEqual(start, end)) {
+            if (constant->AllEqual()) {
                 {
                     ScopedParen sp(out);
-                    if (!EmitConstantRange(out, constant, v->type(), start, start + 1)) {
+                    if (!EmitConstant(out, constant->Index(0))) {
                         return false;
                     }
                 }
                 out << ".";
-                for (size_t i = start; i < end; i++) {
+                for (size_t i = 0; i < v->Width(); i++) {
                     out << "x";
                 }
                 return true;
@@ -3157,11 +3151,11 @@
 
             ScopedParen sp(out);
 
-            for (size_t i = start; i < end; i++) {
-                if (i > start) {
+            for (size_t i = 0; i < v->Width(); i++) {
+                if (i > 0) {
                     out << ", ";
                 }
-                if (!EmitConstantRange(out, constant, v->type(), i, i + 1u)) {
+                if (!EmitConstant(out, constant->Index(i))) {
                     return false;
                 }
             }
@@ -3174,20 +3168,18 @@
 
             ScopedParen sp(out);
 
-            for (size_t column_idx = 0; column_idx < m->columns(); column_idx++) {
-                if (column_idx > 0) {
+            for (size_t i = 0; i < m->columns(); i++) {
+                if (i > 0) {
                     out << ", ";
                 }
-                size_t col_start = m->rows() * column_idx;
-                size_t col_end = col_start + m->rows();
-                if (!EmitConstantRange(out, constant, m->ColumnType(), col_start, col_end)) {
+                if (!EmitConstant(out, constant->Index(i))) {
                     return false;
                 }
             }
             return true;
         },
         [&](const sem::Array* a) {
-            if (constant.AllZero(start, end)) {
+            if (constant->AllZero()) {
                 out << "(";
                 if (!EmitType(out, a, ast::StorageClass::kNone, ast::Access::kUndefined, "")) {
                     return false;
@@ -3199,15 +3191,11 @@
             out << "{";
             TINT_DEFER(out << "}");
 
-            auto* el_ty = a->ElemType();
-
-            uint32_t step = 0;
-            sem::Type::DeepestElementOf(el_ty, &step);
-            for (size_t i = start; i < end; i += step) {
-                if (i > start) {
+            for (size_t i = 0; i < a->Count(); i++) {
+                if (i > 0) {
                     out << ", ";
                 }
-                if (!EmitConstantRange(out, constant, el_ty, i, i + step)) {
+                if (!EmitConstant(out, constant->Index(i))) {
                     return false;
                 }
             }
@@ -3217,7 +3205,7 @@
         [&](Default) {
             diagnostics_.add_error(
                 diag::System::Writer,
-                "unhandled constant type: " + builder_.FriendlyName(constant.Type()));
+                "unhandled constant type: " + builder_.FriendlyName(constant->Type()));
             return false;
         });
 }
diff --git a/src/tint/writer/hlsl/generator_impl.h b/src/tint/writer/hlsl/generator_impl.h
index 8960548..bf2debe 100644
--- a/src/tint/writer/hlsl/generator_impl.h
+++ b/src/tint/writer/hlsl/generator_impl.h
@@ -93,7 +93,8 @@
     /// @param stmt the statement to emit
     /// @returns true if the statement was emitted successfully
     bool EmitAssign(const ast::AssignmentStatement* stmt);
-    /// Emits code such that if `expr` is zero, it emits one, else `expr`
+    /// Emits code such that if `expr` is zero, it emits one, else `expr`.
+    /// Used to avoid divide-by-zeros by substituting constant zeros with ones.
     /// @param out the output of the expression stream
     /// @param expr the expression
     /// @returns true if the expression was emitted, false otherwise
@@ -342,19 +343,7 @@
     /// @param out the output stream
     /// @param constant the constant value to emit
     /// @returns true if the constant value was successfully emitted
-    bool EmitConstant(std::ostream& out, const sem::Constant& constant);
-    /// Handles emitting a sub-range of a constant value
-    /// @param out the output stream
-    /// @param constant the constant value to emit
-    /// @param range_ty the sub-range type
-    /// @param start the element index for the first element
-    /// @param end the element index for one past the last element
-    /// @returns true if the constant value was successfully emitted
-    bool EmitConstantRange(std::ostream& out,
-                           const sem::Constant& constant,
-                           const sem::Type* range_ty,
-                           size_t start,
-                           size_t end);
+    bool EmitConstant(std::ostream& out, const sem::Constant* constant);
     /// Handles a literal
     /// @param out the output stream
     /// @param lit the literal to emit
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index f5ffa6c..43324bb 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -1201,7 +1201,7 @@
                     break;  // Other texture dimensions don't have an offset
             }
         }
-        auto c = component->ConstantValue().Element<AInt>(0);
+        auto c = component->ConstantValue()->As<AInt>();
         switch (c.value) {
             case 0:
                 out << "component::x";
@@ -1594,31 +1594,23 @@
         });
 }
 
-bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant& constant) {
-    return EmitConstantRange(out, constant, constant.Type(), 0, constant.ElementCount());
-}
-
-bool GeneratorImpl::EmitConstantRange(std::ostream& out,
-                                      const sem::Constant& constant,
-                                      const sem::Type* range_ty,
-                                      size_t start,
-                                      size_t end) {
+bool GeneratorImpl::EmitConstant(std::ostream& out, const sem::Constant* constant) {
     return Switch(
-        range_ty,  //
+        constant->Type(),  //
         [&](const sem::Bool*) {
-            out << (constant.Element<AInt>(start) ? "true" : "false");
+            out << (constant->As<AInt>() ? "true" : "false");
             return true;
         },
         [&](const sem::F32*) {
-            PrintF32(out, static_cast<float>(constant.Element<AFloat>(start)));
+            PrintF32(out, constant->As<float>());
             return true;
         },
         [&](const sem::I32*) {
-            PrintI32(out, static_cast<int32_t>(constant.Element<AInt>(start).value));
+            PrintI32(out, constant->As<int32_t>());
             return true;
         },
         [&](const sem::U32*) {
-            out << constant.Element<AInt>(start).value << "u";
+            out << constant->As<AInt>() << "u";
             return true;
         },
         [&](const sem::Vector* v) {
@@ -1628,18 +1620,18 @@
 
             ScopedParen sp(out);
 
-            if (constant.AllEqual(start, end)) {
-                if (!EmitConstantRange(out, constant, v->type(), start, start + 1)) {
+            if (constant->AllEqual()) {
+                if (!EmitConstant(out, constant->Index(0))) {
                     return false;
                 }
                 return true;
             }
 
-            for (size_t i = start; i < end; i++) {
-                if (i > start) {
+            for (size_t i = 0; i < v->Width(); i++) {
+                if (i > 0) {
                     out << ", ";
                 }
-                if (!EmitConstantRange(out, constant, v->type(), i, i + 1u)) {
+                if (!EmitConstant(out, constant->Index(i))) {
                     return false;
                 }
             }
@@ -1652,13 +1644,11 @@
 
             ScopedParen sp(out);
 
-            for (size_t column_idx = 0; column_idx < m->columns(); column_idx++) {
-                if (column_idx > 0) {
+            for (size_t i = 0; i < m->columns(); i++) {
+                if (i > 0) {
                     out << ", ";
                 }
-                size_t col_start = m->rows() * column_idx;
-                size_t col_end = col_start + m->rows();
-                if (!EmitConstantRange(out, constant, m->ColumnType(), col_start, col_end)) {
+                if (!EmitConstant(out, constant->Index(i))) {
                     return false;
                 }
             }
@@ -1669,7 +1659,7 @@
                 return false;
             }
 
-            if (constant.AllZero(start, end)) {
+            if (constant->AllZero()) {
                 out << "{}";
                 return true;
             }
@@ -1677,15 +1667,11 @@
             out << "{";
             TINT_DEFER(out << "}");
 
-            auto* el_ty = a->ElemType();
-
-            uint32_t step = 0;
-            sem::Type::DeepestElementOf(el_ty, &step);
-            for (size_t i = start; i < end; i += step) {
-                if (i > start) {
+            for (size_t i = 0; i < a->Count(); i++) {
+                if (i > 0) {
                     out << ", ";
                 }
-                if (!EmitConstantRange(out, constant, el_ty, i, i + step)) {
+                if (!EmitConstant(out, constant->Index(i))) {
                     return false;
                 }
             }
@@ -1695,7 +1681,7 @@
         [&](Default) {
             diagnostics_.add_error(
                 diag::System::Writer,
-                "unhandled constant type: " + builder_.FriendlyName(constant.Type()));
+                "unhandled constant type: " + builder_.FriendlyName(constant->Type()));
             return false;
         });
 }
diff --git a/src/tint/writer/msl/generator_impl.h b/src/tint/writer/msl/generator_impl.h
index ed0c35b..f27d258 100644
--- a/src/tint/writer/msl/generator_impl.h
+++ b/src/tint/writer/msl/generator_impl.h
@@ -256,19 +256,7 @@
     /// @param out the output stream
     /// @param constant the constant value to emit
     /// @returns true if the constant value was successfully emitted
-    bool EmitConstant(std::ostream& out, const sem::Constant& constant);
-    /// Handles emitting a sub-range of a constant value
-    /// @param out the output stream
-    /// @param constant the constant value to emit
-    /// @param range_ty the sub-range type
-    /// @param start the element index for the first element
-    /// @param end the element index for one past the last element
-    /// @returns true if the constant value was successfully emitted
-    bool EmitConstantRange(std::ostream& out,
-                           const sem::Constant& constant,
-                           const sem::Type* range_ty,
-                           size_t start,
-                           size_t end);
+    bool EmitConstant(std::ostream& out, const sem::Constant* constant);
     /// Handles a literal
     /// @param out the output of the expression stream
     /// @param lit the literal to emit
diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc
index fea7a14..833324a 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -959,7 +959,7 @@
                                     Operand(result_type_id),
                                     extract,
                                     Operand(info->source_id),
-                                    Operand(idx_constval.Element<uint32_t>(0)),
+                                    Operand(idx_constval->As<uint32_t>()),
                                 })) {
             return false;
         }
@@ -1703,20 +1703,14 @@
     return GenerateConstantIfNeeded(constant);
 }
 
-uint32_t Builder::GenerateConstantIfNeeded(const sem::Constant& constant) {
-    return GenerateConstantRangeIfNeeded(constant, constant.Type(), 0, constant.ElementCount());
-}
-
-uint32_t Builder::GenerateConstantRangeIfNeeded(const sem::Constant& constant,
-                                                const sem::Type* range_ty,
-                                                size_t start,
-                                                size_t end) {
-    if (constant.AllZero(start, end)) {
-        return GenerateConstantNullIfNeeded(range_ty);
+uint32_t Builder::GenerateConstantIfNeeded(const sem::Constant* constant) {
+    if (constant->AllZero()) {
+        return GenerateConstantNullIfNeeded(constant->Type());
     }
+    auto* ty = constant->Type();
 
-    auto composite = [&](const sem::Type* el_ty) -> uint32_t {
-        auto type_id = GenerateTypeIfNeeded(range_ty);
+    auto composite = [&](size_t el_count) -> uint32_t {
+        auto type_id = GenerateTypeIfNeeded(ty);
         if (!type_id) {
             return 0;
         }
@@ -1724,14 +1718,12 @@
         static constexpr size_t kOpsResultIdx = 1;  // operand index of the result
 
         std::vector<Operand> ops;
-        ops.reserve(end - start + 2);
+        ops.reserve(el_count + 2);
         ops.emplace_back(type_id);
         ops.push_back(Operand(0u));  // Placeholder for the result ID
 
-        uint32_t step = 0;
-        sem::Type::DeepestElementOf(el_ty, &step);
-        for (size_t i = start; i < end; i += step) {
-            auto id = GenerateConstantRangeIfNeeded(constant, el_ty, i, i + step);
+        for (size_t i = 0; i < el_count; i++) {
+            auto id = GenerateConstantIfNeeded(constant->Index(i));
             if (!id) {
                 return 0;
             }
@@ -1749,28 +1741,28 @@
     };
 
     return Switch(
-        range_ty,  //
+        ty,  //
         [&](const sem::Bool*) {
-            bool val = constant.Element<AInt>(start);
+            bool val = constant->As<bool>();
             return GenerateConstantIfNeeded(ScalarConstant::Bool(val));
         },
         [&](const sem::F32*) {
-            auto val = f32(constant.Element<AFloat>(start));
+            auto val = constant->As<f32>();
             return GenerateConstantIfNeeded(ScalarConstant::F32(val.value));
         },
         [&](const sem::I32*) {
-            auto val = i32(constant.Element<AInt>(start));
+            auto val = constant->As<i32>();
             return GenerateConstantIfNeeded(ScalarConstant::I32(val.value));
         },
         [&](const sem::U32*) {
-            auto val = u32(constant.Element<AInt>(start));
+            auto val = constant->As<u32>();
             return GenerateConstantIfNeeded(ScalarConstant::U32(val.value));
         },
-        [&](const sem::Vector* v) { return composite(v->type()); },
-        [&](const sem::Matrix* m) { return composite(m->ColumnType()); },
-        [&](const sem::Array* a) { return composite(a->ElemType()); },
+        [&](const sem::Vector* v) { return composite(v->Width()); },
+        [&](const sem::Matrix* m) { return composite(m->columns()); },
+        [&](const sem::Array* a) { return composite(a->Count()); },
         [&](Default) {
-            error_ = "unhandled constant type: " + builder_.FriendlyName(constant.Type());
+            error_ = "unhandled constant type: " + builder_.FriendlyName(ty);
             return false;
         });
 }
diff --git a/src/tint/writer/spirv/builder.h b/src/tint/writer/spirv/builder.h
index 8465089..f6bb93f 100644
--- a/src/tint/writer/spirv/builder.h
+++ b/src/tint/writer/spirv/builder.h
@@ -554,18 +554,7 @@
     /// Generates a constant value if needed
     /// @param constant the constant to generate.
     /// @returns the ID on success or 0 on failure
-    uint32_t GenerateConstantIfNeeded(const sem::Constant& constant);
-
-    /// Handles emitting a sub-range of a constant value
-    /// @param constant the constant value to emit
-    /// @param range_ty the sub-range type
-    /// @param start the element index for the first element
-    /// @param end the element index for one past the last element
-    /// @returns true if the constant value was successfully emitted
-    uint32_t GenerateConstantRangeIfNeeded(const sem::Constant& constant,
-                                           const sem::Type* range_ty,
-                                           size_t start,
-                                           size_t end);
+    uint32_t GenerateConstantIfNeeded(const sem::Constant* constant);
 
     /// Generates a scalar constant if needed
     /// @param constant the constant to generate.