sem: Use UniqueAllocator for types

Add a Hash() and Equals() methods to sem::Type.

Have sem::Manager (which should be called sem::TypeManager) derive from
utils::UniqueAllocator. This now uses the Hash() and Equals() for uniquely
constructing semantic types instead of building strings and comparing those.

Bug: tint:1383
Change-Id: I5e3229bd087391ac594d333a0ab4232cfcddf54d
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/82743
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
diff --git a/src/tint/builtin_table.cc b/src/tint/builtin_table.cc
index 4d11c9d..05c5673 100644
--- a/src/tint/builtin_table.cc
+++ b/src/tint/builtin_table.cc
@@ -47,6 +47,11 @@
  public:
   Any() = default;
   ~Any() override = default;
+
+  // Stub implementations for sem::Type conformance.
+  size_t Hash() const override { return 0; }
+  bool Equals(const sem::Type&) const override { return false; }
+
   std::string type_name() const override { return "<any>"; }
   std::string FriendlyName(const SymbolTable&) const override {
     return "<any>";
@@ -429,11 +434,11 @@
 
 const sem::Array* build_array(MatchState& state, const sem::Type* el) {
   return state.builder.create<sem::Array>(el,
-                                          /* count */ 0,
-                                          /* align */ 0,
-                                          /* size */ 0,
-                                          /* stride */ 0,
-                                          /* stride_implicit */ 0);
+                                          /* count */ 0u,
+                                          /* align */ 0u,
+                                          /* size */ 0u,
+                                          /* stride */ 0u,
+                                          /* stride_implicit */ 0u);
 }
 
 bool match_ptr(const sem::Type* ty, Number& S, const sem::Type*& T, Number& A) {
diff --git a/src/tint/builtin_table_test.cc b/src/tint/builtin_table_test.cc
index 1734c7d..b8454d1 100644
--- a/src/tint/builtin_table_test.cc
+++ b/src/tint/builtin_table_test.cc
@@ -60,7 +60,7 @@
 TEST_F(BuiltinTableTest, MatchU32) {
   auto* f32 = create<sem::F32>();
   auto* u32 = create<sem::U32>();
-  auto* vec2_f32 = create<sem::Vector>(f32, 2);
+  auto* vec2_f32 = create<sem::Vector>(f32, 2u);
   auto* result = table->Lookup(BuiltinType::kUnpack2x16float, {u32}, Source{});
   ASSERT_NE(result, nullptr) << Diagnostics().str();
   ASSERT_EQ(Diagnostics().str(), "");
@@ -80,7 +80,7 @@
 TEST_F(BuiltinTableTest, MatchI32) {
   auto* f32 = create<sem::F32>();
   auto* i32 = create<sem::I32>();
-  auto* vec4_f32 = create<sem::Vector>(f32, 4);
+  auto* vec4_f32 = create<sem::Vector>(f32, 4u);
   auto* tex = create<sem::SampledTexture>(ast::TextureDimension::k1d, f32);
   auto* result =
       table->Lookup(BuiltinType::kTextureLoad, {tex, i32, i32}, Source{});
@@ -226,7 +226,7 @@
 }
 
 TEST_F(BuiltinTableTest, MatchArray) {
-  auto* arr = create<sem::Array>(create<sem::U32>(), 0, 4, 4, 4, 4);
+  auto* arr = create<sem::Array>(create<sem::U32>(), 0u, 4u, 4u, 4u, 4u);
   auto* arr_ptr = create<sem::Pointer>(arr, ast::StorageClass::kStorage,
                                        ast::Access::kReadWrite);
   auto* result = table->Lookup(BuiltinType::kArrayLength, {arr_ptr}, Source{});
@@ -249,8 +249,8 @@
 
 TEST_F(BuiltinTableTest, MatchSampler) {
   auto* f32 = create<sem::F32>();
-  auto* vec2_f32 = create<sem::Vector>(f32, 2);
-  auto* vec4_f32 = create<sem::Vector>(f32, 4);
+  auto* vec2_f32 = create<sem::Vector>(f32, 2u);
+  auto* vec4_f32 = create<sem::Vector>(f32, 4u);
   auto* tex = create<sem::SampledTexture>(ast::TextureDimension::k2d, f32);
   auto* sampler = create<sem::Sampler>(ast::SamplerKind::kSampler);
   auto* result = table->Lookup(BuiltinType::kTextureSample,
@@ -270,7 +270,7 @@
 
 TEST_F(BuiltinTableTest, MismatchSampler) {
   auto* f32 = create<sem::F32>();
-  auto* vec2_f32 = create<sem::Vector>(f32, 2);
+  auto* vec2_f32 = create<sem::Vector>(f32, 2u);
   auto* tex = create<sem::SampledTexture>(ast::TextureDimension::k2d, f32);
   auto* result = table->Lookup(BuiltinType::kTextureSample,
                                {tex, f32, vec2_f32}, Source{});
@@ -281,8 +281,8 @@
 TEST_F(BuiltinTableTest, MatchSampledTexture) {
   auto* i32 = create<sem::I32>();
   auto* f32 = create<sem::F32>();
-  auto* vec2_i32 = create<sem::Vector>(i32, 2);
-  auto* vec4_f32 = create<sem::Vector>(f32, 4);
+  auto* vec2_i32 = create<sem::Vector>(i32, 2u);
+  auto* vec4_f32 = create<sem::Vector>(f32, 4u);
   auto* tex = create<sem::SampledTexture>(ast::TextureDimension::k2d, f32);
   auto* result =
       table->Lookup(BuiltinType::kTextureLoad, {tex, vec2_i32, i32}, Source{});
@@ -302,8 +302,8 @@
 TEST_F(BuiltinTableTest, MatchMultisampledTexture) {
   auto* i32 = create<sem::I32>();
   auto* f32 = create<sem::F32>();
-  auto* vec2_i32 = create<sem::Vector>(i32, 2);
-  auto* vec4_f32 = create<sem::Vector>(f32, 4);
+  auto* vec2_i32 = create<sem::Vector>(i32, 2u);
+  auto* vec4_f32 = create<sem::Vector>(f32, 4u);
   auto* tex = create<sem::MultisampledTexture>(ast::TextureDimension::k2d, f32);
   auto* result =
       table->Lookup(BuiltinType::kTextureLoad, {tex, vec2_i32, i32}, Source{});
@@ -323,7 +323,7 @@
 TEST_F(BuiltinTableTest, MatchDepthTexture) {
   auto* f32 = create<sem::F32>();
   auto* i32 = create<sem::I32>();
-  auto* vec2_i32 = create<sem::Vector>(i32, 2);
+  auto* vec2_i32 = create<sem::Vector>(i32, 2u);
   auto* tex = create<sem::DepthTexture>(ast::TextureDimension::k2d);
   auto* result =
       table->Lookup(BuiltinType::kTextureLoad, {tex, vec2_i32, i32}, Source{});
@@ -343,7 +343,7 @@
 TEST_F(BuiltinTableTest, MatchDepthMultisampledTexture) {
   auto* f32 = create<sem::F32>();
   auto* i32 = create<sem::I32>();
-  auto* vec2_i32 = create<sem::Vector>(i32, 2);
+  auto* vec2_i32 = create<sem::Vector>(i32, 2u);
   auto* tex = create<sem::DepthMultisampledTexture>(ast::TextureDimension::k2d);
   auto* result =
       table->Lookup(BuiltinType::kTextureLoad, {tex, vec2_i32, i32}, Source{});
@@ -363,8 +363,8 @@
 TEST_F(BuiltinTableTest, MatchExternalTexture) {
   auto* f32 = create<sem::F32>();
   auto* i32 = create<sem::I32>();
-  auto* vec2_i32 = create<sem::Vector>(i32, 2);
-  auto* vec4_f32 = create<sem::Vector>(f32, 4);
+  auto* vec2_i32 = create<sem::Vector>(i32, 2u);
+  auto* vec4_f32 = create<sem::Vector>(f32, 4u);
   auto* tex = create<sem::ExternalTexture>();
   auto* result =
       table->Lookup(BuiltinType::kTextureLoad, {tex, vec2_i32}, Source{});
@@ -382,8 +382,8 @@
 TEST_F(BuiltinTableTest, MatchWOStorageTexture) {
   auto* f32 = create<sem::F32>();
   auto* i32 = create<sem::I32>();
-  auto* vec2_i32 = create<sem::Vector>(i32, 2);
-  auto* vec4_f32 = create<sem::Vector>(f32, 4);
+  auto* vec2_i32 = create<sem::Vector>(i32, 2u);
+  auto* vec4_f32 = create<sem::Vector>(f32, 4u);
   auto* subtype =
       sem::StorageTexture::SubtypeFor(ast::TexelFormat::kR32Float, Types());
   auto* tex = create<sem::StorageTexture>(ast::TextureDimension::k2d,
@@ -408,7 +408,7 @@
 TEST_F(BuiltinTableTest, MismatchTexture) {
   auto* f32 = create<sem::F32>();
   auto* i32 = create<sem::I32>();
-  auto* vec2_i32 = create<sem::Vector>(i32, 2);
+  auto* vec2_i32 = create<sem::Vector>(i32, 2u);
   auto* result =
       table->Lookup(BuiltinType::kTextureLoad, {f32, vec2_i32}, Source{});
   ASSERT_EQ(result, nullptr);
@@ -452,7 +452,7 @@
 
 TEST_F(BuiltinTableTest, MatchOpenSizeVector) {
   auto* f32 = create<sem::F32>();
-  auto* vec2_f32 = create<sem::Vector>(f32, 2);
+  auto* vec2_f32 = create<sem::Vector>(f32, 2u);
   auto* result = table->Lookup(BuiltinType::kClamp,
                                {vec2_f32, vec2_f32, vec2_f32}, Source{});
   ASSERT_NE(result, nullptr) << Diagnostics().str();
@@ -468,7 +468,7 @@
 TEST_F(BuiltinTableTest, MismatchOpenSizeVector) {
   auto* f32 = create<sem::F32>();
   auto* u32 = create<sem::U32>();
-  auto* vec2_f32 = create<sem::Vector>(f32, 2);
+  auto* vec2_f32 = create<sem::Vector>(f32, 2u);
   auto* result =
       table->Lookup(BuiltinType::kClamp, {vec2_f32, u32, vec2_f32}, Source{});
   ASSERT_EQ(result, nullptr);
@@ -477,8 +477,8 @@
 
 TEST_F(BuiltinTableTest, MatchOpenSizeMatrix) {
   auto* f32 = create<sem::F32>();
-  auto* vec3_f32 = create<sem::Vector>(f32, 3);
-  auto* mat3_f32 = create<sem::Matrix>(vec3_f32, 3);
+  auto* vec3_f32 = create<sem::Vector>(f32, 3u);
+  auto* mat3_f32 = create<sem::Matrix>(vec3_f32, 3u);
   auto* result = table->Lookup(BuiltinType::kDeterminant, {mat3_f32}, Source{});
   ASSERT_NE(result, nullptr) << Diagnostics().str();
   ASSERT_EQ(Diagnostics().str(), "");
@@ -490,8 +490,8 @@
 
 TEST_F(BuiltinTableTest, MismatchOpenSizeMatrix) {
   auto* f32 = create<sem::F32>();
-  auto* vec2_f32 = create<sem::Vector>(f32, 2);
-  auto* mat3x2_f32 = create<sem::Matrix>(vec2_f32, 3);
+  auto* vec2_f32 = create<sem::Vector>(f32, 2u);
+  auto* mat3x2_f32 = create<sem::Matrix>(vec2_f32, 3u);
   auto* result =
       table->Lookup(BuiltinType::kDeterminant, {mat3x2_f32}, Source{});
   ASSERT_EQ(result, nullptr);
@@ -578,7 +578,7 @@
 
 TEST_F(BuiltinTableTest, SameOverloadReturnsSameBuiltinPointer) {
   auto* f32 = create<sem::F32>();
-  auto* vec2_f32 = create<sem::Vector>(create<sem::F32>(), 2);
+  auto* vec2_f32 = create<sem::Vector>(create<sem::F32>(), 2u);
   auto* bool_ = create<sem::Bool>();
   auto* a = table->Lookup(BuiltinType::kSelect, {f32, f32, bool_}, Source{});
   ASSERT_NE(a, nullptr) << Diagnostics().str();
diff --git a/src/tint/resolver/inferred_type_test.cc b/src/tint/resolver/inferred_type_test.cc
index 689f813..534cb07 100644
--- a/src/tint/resolver/inferred_type_test.cc
+++ b/src/tint/resolver/inferred_type_test.cc
@@ -142,7 +142,7 @@
 TEST_F(ResolverInferredTypeTest, InferArray_Pass) {
   auto* type = ty.array(ty.u32(), 10);
   auto* expected_type =
-      create<sem::Array>(create<sem::U32>(), 10, 4, 4 * 10, 4, 4);
+      create<sem::Array>(create<sem::U32>(), 10u, 4u, 4u * 10u, 4u, 4u);
 
   auto* ctor_expr = Construct(type);
   auto* var = Var("a", nullptr, ast::StorageClass::kFunction, ctor_expr);
@@ -159,8 +159,8 @@
   auto* expected_type = create<sem::Struct>(
       str, str->name,
       sem::StructMemberList{create<sem::StructMember>(
-          member, member->symbol, create<sem::I32>(), 0, 0, 0, 4)},
-      0, 4, 4);
+          member, member->symbol, create<sem::I32>(), 0u, 0u, 0u, 4u)},
+      0u, 4u, 4u);
 
   auto* ctor_expr = Construct(ty.Of(str));
 
diff --git a/src/tint/resolver/is_host_shareable_test.cc b/src/tint/resolver/is_host_shareable_test.cc
index 9c992ab..c7069d7 100644
--- a/src/tint/resolver/is_host_shareable_test.cc
+++ b/src/tint/resolver/is_host_shareable_test.cc
@@ -39,52 +39,61 @@
 }
 
 TEST_F(ResolverIsHostShareable, NumericVector) {
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::I32>(), 2)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::I32>(), 3)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::I32>(), 4)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::U32>(), 2)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::U32>(), 3)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::U32>(), 4)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::F32>(), 2)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::F32>(), 3)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Vector>(create<sem::F32>(), 4)));
+  EXPECT_TRUE(
+      r()->IsHostShareable(create<sem::Vector>(create<sem::I32>(), 2u)));
+  EXPECT_TRUE(
+      r()->IsHostShareable(create<sem::Vector>(create<sem::I32>(), 3u)));
+  EXPECT_TRUE(
+      r()->IsHostShareable(create<sem::Vector>(create<sem::I32>(), 4u)));
+  EXPECT_TRUE(
+      r()->IsHostShareable(create<sem::Vector>(create<sem::U32>(), 2u)));
+  EXPECT_TRUE(
+      r()->IsHostShareable(create<sem::Vector>(create<sem::U32>(), 3u)));
+  EXPECT_TRUE(
+      r()->IsHostShareable(create<sem::Vector>(create<sem::U32>(), 4u)));
+  EXPECT_TRUE(
+      r()->IsHostShareable(create<sem::Vector>(create<sem::F32>(), 2u)));
+  EXPECT_TRUE(
+      r()->IsHostShareable(create<sem::Vector>(create<sem::F32>(), 3u)));
+  EXPECT_TRUE(
+      r()->IsHostShareable(create<sem::Vector>(create<sem::F32>(), 4u)));
 }
 
 TEST_F(ResolverIsHostShareable, BoolVector) {
   EXPECT_FALSE(
-      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 2)));
+      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 2u)));
   EXPECT_FALSE(
-      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 3)));
+      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 3u)));
   EXPECT_FALSE(
-      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 4)));
+      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 4u)));
   EXPECT_FALSE(
-      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 2)));
+      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 2u)));
   EXPECT_FALSE(
-      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 3)));
+      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 3u)));
   EXPECT_FALSE(
-      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 4)));
+      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 4u)));
   EXPECT_FALSE(
-      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 2)));
+      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 2u)));
   EXPECT_FALSE(
-      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 3)));
+      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 3u)));
   EXPECT_FALSE(
-      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 4)));
+      r()->IsHostShareable(create<sem::Vector>(create<sem::Bool>(), 4u)));
 }
 
 TEST_F(ResolverIsHostShareable, Matrix) {
-  auto* vec2 = create<sem::Vector>(create<sem::F32>(), 2);
-  auto* vec3 = create<sem::Vector>(create<sem::F32>(), 3);
-  auto* vec4 = create<sem::Vector>(create<sem::F32>(), 4);
+  auto* vec2 = create<sem::Vector>(create<sem::F32>(), 2u);
+  auto* vec3 = create<sem::Vector>(create<sem::F32>(), 3u);
+  auto* vec4 = create<sem::Vector>(create<sem::F32>(), 4u);
 
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2, 2)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2, 3)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2, 4)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3, 2)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3, 3)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3, 4)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4, 2)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4, 3)));
-  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4, 4)));
+  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2, 2u)));
+  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2, 3u)));
+  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec2, 4u)));
+  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3, 2u)));
+  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3, 3u)));
+  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec3, 4u)));
+  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4, 2u)));
+  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4, 3u)));
+  EXPECT_TRUE(r()->IsHostShareable(create<sem::Matrix>(vec4, 4u)));
 }
 
 TEST_F(ResolverIsHostShareable, Pointer) {
@@ -99,12 +108,12 @@
 }
 
 TEST_F(ResolverIsHostShareable, ArraySizedOfHostShareable) {
-  auto* arr = create<sem::Array>(create<sem::I32>(), 5, 4, 20, 4, 4);
+  auto* arr = create<sem::Array>(create<sem::I32>(), 5u, 4u, 20u, 4u, 4u);
   EXPECT_TRUE(r()->IsHostShareable(arr));
 }
 
 TEST_F(ResolverIsHostShareable, ArrayUnsizedOfHostShareable) {
-  auto* arr = create<sem::Array>(create<sem::I32>(), 0, 4, 4, 4, 4);
+  auto* arr = create<sem::Array>(create<sem::I32>(), 0u, 4u, 4u, 4u, 4u);
   EXPECT_TRUE(r()->IsHostShareable(arr));
 }
 
diff --git a/src/tint/resolver/is_storeable_test.cc b/src/tint/resolver/is_storeable_test.cc
index cc3323a..739ae22 100644
--- a/src/tint/resolver/is_storeable_test.cc
+++ b/src/tint/resolver/is_storeable_test.cc
@@ -36,30 +36,30 @@
 }
 
 TEST_F(ResolverIsStorableTest, Vector) {
-  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::I32>(), 2)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::I32>(), 3)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::I32>(), 4)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::U32>(), 2)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::U32>(), 3)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::U32>(), 4)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::F32>(), 2)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::F32>(), 3)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::F32>(), 4)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::I32>(), 2u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::I32>(), 3u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::I32>(), 4u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::U32>(), 2u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::U32>(), 3u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::U32>(), 4u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::F32>(), 2u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::F32>(), 3u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Vector>(create<sem::F32>(), 4u)));
 }
 
 TEST_F(ResolverIsStorableTest, Matrix) {
-  auto* vec2 = create<sem::Vector>(create<sem::F32>(), 2);
-  auto* vec3 = create<sem::Vector>(create<sem::F32>(), 3);
-  auto* vec4 = create<sem::Vector>(create<sem::F32>(), 4);
-  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2, 2)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2, 3)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2, 4)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3, 2)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3, 3)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3, 4)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4, 2)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4, 3)));
-  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4, 4)));
+  auto* vec2 = create<sem::Vector>(create<sem::F32>(), 2u);
+  auto* vec3 = create<sem::Vector>(create<sem::F32>(), 3u);
+  auto* vec4 = create<sem::Vector>(create<sem::F32>(), 4u);
+  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2, 2u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2, 3u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec2, 4u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3, 2u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3, 3u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec3, 4u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4, 2u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4, 3u)));
+  EXPECT_TRUE(r()->IsStorable(create<sem::Matrix>(vec4, 4u)));
 }
 
 TEST_F(ResolverIsStorableTest, Pointer) {
@@ -74,12 +74,12 @@
 }
 
 TEST_F(ResolverIsStorableTest, ArraySizedOfStorable) {
-  auto* arr = create<sem::Array>(create<sem::I32>(), 5, 4, 20, 4, 4);
+  auto* arr = create<sem::Array>(create<sem::I32>(), 5u, 4u, 20u, 4u, 4u);
   EXPECT_TRUE(r()->IsStorable(arr));
 }
 
 TEST_F(ResolverIsStorableTest, ArrayUnsizedOfStorable) {
-  auto* arr = create<sem::Array>(create<sem::I32>(), 0, 4, 4, 4, 4);
+  auto* arr = create<sem::Array>(create<sem::I32>(), 0u, 4u, 4u, 4u, 4u);
   EXPECT_TRUE(r()->IsStorable(arr));
 }
 
diff --git a/src/tint/resolver/resolver_test_helper.h b/src/tint/resolver/resolver_test_helper.h
index b128a77..74c7eb3 100644
--- a/src/tint/resolver/resolver_test_helper.h
+++ b/src/tint/resolver/resolver_test_helper.h
@@ -126,7 +126,7 @@
 using u32 = ProgramBuilder::u32;
 using f32 = ProgramBuilder::f32;
 
-template <int N, typename T>
+template <uint32_t N, typename T>
 struct vec {};
 
 template <typename T>
@@ -138,7 +138,7 @@
 template <typename T>
 using vec4 = vec<4, T>;
 
-template <int N, int M, typename T>
+template <uint32_t N, uint32_t M, typename T>
 struct mat {};
 
 template <typename T>
@@ -156,7 +156,7 @@
 template <typename T>
 using mat4x4 = mat<4, 4, T>;
 
-template <int N, typename T>
+template <uint32_t N, typename T>
 struct array {};
 
 template <typename TO, int ID = 0>
@@ -271,7 +271,7 @@
 };
 
 /// Helper for building vector types and expressions
-template <int N, typename T>
+template <uint32_t N, typename T>
 struct DataType<vec<N, T>> {
   /// true as vectors are a composite type
   static constexpr bool is_composite = true;
@@ -300,7 +300,7 @@
   static inline ast::ExpressionList ExprArgs(ProgramBuilder& b,
                                              int elem_value) {
     ast::ExpressionList args;
-    for (int i = 0; i < N; i++) {
+    for (uint32_t i = 0; i < N; i++) {
       args.emplace_back(DataType<T>::Expr(b, elem_value));
     }
     return args;
@@ -308,7 +308,7 @@
 };
 
 /// Helper for building matrix types and expressions
-template <int N, int M, typename T>
+template <uint32_t N, uint32_t M, typename T>
 struct DataType<mat<N, M, T>> {
   /// true as matrices are a composite type
   static constexpr bool is_composite = true;
@@ -338,7 +338,7 @@
   static inline ast::ExpressionList ExprArgs(ProgramBuilder& b,
                                              int elem_value) {
     ast::ExpressionList args;
-    for (int i = 0; i < N; i++) {
+    for (uint32_t i = 0; i < N; i++) {
       args.emplace_back(DataType<vec<M, T>>::Expr(b, elem_value));
     }
     return args;
@@ -421,7 +421,7 @@
 };
 
 /// Helper for building array types and expressions
-template <int N, typename T>
+template <uint32_t N, typename T>
 struct DataType<array<N, T>> {
   /// true as arrays are a composite type
   static constexpr bool is_composite = true;
@@ -457,7 +457,7 @@
   static inline ast::ExpressionList ExprArgs(ProgramBuilder& b,
                                              int elem_value) {
     ast::ExpressionList args;
-    for (int i = 0; i < N; i++) {
+    for (uint32_t i = 0; i < N; i++) {
       args.emplace_back(DataType<T>::Expr(b, elem_value));
     }
     return args;
diff --git a/src/tint/sem/array.cc b/src/tint/sem/array.cc
index 8b20aa6..ca79701 100644
--- a/src/tint/sem/array.cc
+++ b/src/tint/sem/array.cc
@@ -17,6 +17,7 @@
 #include <string>
 
 #include "src/tint/debug.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::Array);
 
@@ -40,6 +41,21 @@
   TINT_ASSERT(Semantic, element_);
 }
 
+size_t Array::Hash() const {
+  return utils::Hash(TypeInfo::Of<Array>().full_hashcode, count_, align_, size_,
+                     stride_);
+}
+
+bool Array::Equals(const sem::Type& other) const {
+  if (auto* o = other.As<Array>()) {
+    // Note: implicit_stride is not part of the type_name string as this is
+    // derived from the element type
+    return o->element_ == element_ && o->count_ == count_ &&
+           o->align_ == align_ && o->size_ == size_ && o->stride_ == stride_;
+  }
+  return false;
+}
+
 bool Array::IsConstructible() const {
   return constructible_;
 }
diff --git a/src/tint/sem/array.h b/src/tint/sem/array.h
index bba467d..4ce1c23 100644
--- a/src/tint/sem/array.h
+++ b/src/tint/sem/array.h
@@ -52,6 +52,13 @@
         uint32_t stride,
         uint32_t implicit_stride);
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @return the array element type
   Type const* ElemType() const { return element_; }
 
diff --git a/src/tint/sem/atomic_type.cc b/src/tint/sem/atomic_type.cc
index a57639b..df3e4c7 100644
--- a/src/tint/sem/atomic_type.cc
+++ b/src/tint/sem/atomic_type.cc
@@ -16,6 +16,7 @@
 
 #include "src/tint/program_builder.h"
 #include "src/tint/sem/reference_type.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::Atomic);
 
@@ -26,6 +27,17 @@
   TINT_ASSERT(AST, !subtype->Is<Reference>());
 }
 
+size_t Atomic::Hash() const {
+  return utils::Hash(TypeInfo::Of<Atomic>().full_hashcode, subtype_);
+}
+
+bool Atomic::Equals(const sem::Type& other) const {
+  if (auto* o = other.As<Atomic>()) {
+    return o->subtype_ == subtype_;
+  }
+  return false;
+}
+
 std::string Atomic::type_name() const {
   std::ostringstream out;
   out << "__atomic" << subtype_->type_name();
diff --git a/src/tint/sem/atomic_type.h b/src/tint/sem/atomic_type.h
index 7a70d0a..f7138e7 100644
--- a/src/tint/sem/atomic_type.h
+++ b/src/tint/sem/atomic_type.h
@@ -33,6 +33,13 @@
   Atomic(Atomic&&);
   ~Atomic() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the atomic type
   const sem::Type* Type() const { return subtype_; }
 
diff --git a/src/tint/sem/bool_type.cc b/src/tint/sem/bool_type.cc
index abbc23c..e7a7943 100644
--- a/src/tint/sem/bool_type.cc
+++ b/src/tint/sem/bool_type.cc
@@ -27,6 +27,14 @@
 
 Bool::~Bool() = default;
 
+size_t Bool::Hash() const {
+  return TypeInfo::Of<Bool>().full_hashcode;
+}
+
+bool Bool::Equals(const Type& other) const {
+  return other.Is<Bool>();
+}
+
 std::string Bool::type_name() const {
   return "__bool";
 }
diff --git a/src/tint/sem/bool_type.h b/src/tint/sem/bool_type.h
index 85a8977..0562f06 100644
--- a/src/tint/sem/bool_type.h
+++ b/src/tint/sem/bool_type.h
@@ -37,6 +37,13 @@
   Bool(Bool&&);
   ~Bool() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the name for this type
   std::string type_name() const override;
 
diff --git a/src/tint/sem/depth_multisampled_texture_type.cc b/src/tint/sem/depth_multisampled_texture_type.cc
index 3130c9a..08e824c 100644
--- a/src/tint/sem/depth_multisampled_texture_type.cc
+++ b/src/tint/sem/depth_multisampled_texture_type.cc
@@ -15,6 +15,7 @@
 #include "src/tint/sem/depth_multisampled_texture_type.h"
 
 #include "src/tint/program_builder.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::DepthMultisampledTexture);
 
@@ -38,6 +39,18 @@
 
 DepthMultisampledTexture::~DepthMultisampledTexture() = default;
 
+size_t DepthMultisampledTexture::Hash() const {
+  return utils::Hash(TypeInfo::Of<DepthMultisampledTexture>().full_hashcode,
+                     dim());
+}
+
+bool DepthMultisampledTexture::Equals(const sem::Type& other) const {
+  if (auto* o = other.As<DepthMultisampledTexture>()) {
+    return o->dim() == dim();
+  }
+  return false;
+}
+
 std::string DepthMultisampledTexture::type_name() const {
   std::ostringstream out;
   out << "__depth_multisampled_texture_" << dim();
diff --git a/src/tint/sem/depth_multisampled_texture_type.h b/src/tint/sem/depth_multisampled_texture_type.h
index 08b1239..da4b18a 100644
--- a/src/tint/sem/depth_multisampled_texture_type.h
+++ b/src/tint/sem/depth_multisampled_texture_type.h
@@ -33,6 +33,13 @@
   DepthMultisampledTexture(DepthMultisampledTexture&&);
   ~DepthMultisampledTexture() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the name for this type
   std::string type_name() const override;
 
diff --git a/src/tint/sem/depth_texture_type.cc b/src/tint/sem/depth_texture_type.cc
index 751d042..67bb008 100644
--- a/src/tint/sem/depth_texture_type.cc
+++ b/src/tint/sem/depth_texture_type.cc
@@ -15,6 +15,7 @@
 #include "src/tint/sem/depth_texture_type.h"
 
 #include "src/tint/program_builder.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::DepthTexture);
 
@@ -39,6 +40,17 @@
 
 DepthTexture::~DepthTexture() = default;
 
+size_t DepthTexture::Hash() const {
+  return utils::Hash(TypeInfo::Of<DepthTexture>().full_hashcode, dim());
+}
+
+bool DepthTexture::Equals(const sem::Type& other) const {
+  if (auto* o = other.As<DepthTexture>()) {
+    return o->dim() == dim();
+  }
+  return false;
+}
+
 std::string DepthTexture::type_name() const {
   std::ostringstream out;
   out << "__depth_texture_" << dim();
diff --git a/src/tint/sem/depth_texture_type.h b/src/tint/sem/depth_texture_type.h
index a8315ed..312b37f 100644
--- a/src/tint/sem/depth_texture_type.h
+++ b/src/tint/sem/depth_texture_type.h
@@ -32,6 +32,13 @@
   DepthTexture(DepthTexture&&);
   ~DepthTexture() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the name for this type
   std::string type_name() const override;
 
diff --git a/src/tint/sem/external_texture_type.cc b/src/tint/sem/external_texture_type.cc
index 5e71d90..f3cf01a 100644
--- a/src/tint/sem/external_texture_type.cc
+++ b/src/tint/sem/external_texture_type.cc
@@ -27,6 +27,14 @@
 
 ExternalTexture::~ExternalTexture() = default;
 
+size_t ExternalTexture::Hash() const {
+  return TypeInfo::Of<ExternalTexture>().full_hashcode;
+}
+
+bool ExternalTexture::Equals(const sem::Type& other) const {
+  return other.Is<ExternalTexture>();
+}
+
 std::string ExternalTexture::type_name() const {
   return "__external_texture";
 }
diff --git a/src/tint/sem/external_texture_type.h b/src/tint/sem/external_texture_type.h
index 6e383a8..5f75303 100644
--- a/src/tint/sem/external_texture_type.h
+++ b/src/tint/sem/external_texture_type.h
@@ -32,6 +32,13 @@
   ExternalTexture(ExternalTexture&&);
   ~ExternalTexture() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the name for this type
   std::string type_name() const override;
 
diff --git a/src/tint/sem/f32_type.cc b/src/tint/sem/f32_type.cc
index 7a071ee..7bdd99e 100644
--- a/src/tint/sem/f32_type.cc
+++ b/src/tint/sem/f32_type.cc
@@ -27,6 +27,14 @@
 
 F32::~F32() = default;
 
+size_t F32::Hash() const {
+  return TypeInfo::Of<F32>().full_hashcode;
+}
+
+bool F32::Equals(const Type& other) const {
+  return other.Is<F32>();
+}
+
 std::string F32::type_name() const {
   return "__f32";
 }
diff --git a/src/tint/sem/f32_type.h b/src/tint/sem/f32_type.h
index ddca10f..ee01d49 100644
--- a/src/tint/sem/f32_type.h
+++ b/src/tint/sem/f32_type.h
@@ -31,6 +31,13 @@
   F32(F32&&);
   ~F32() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the name for this type
   std::string type_name() const override;
 
diff --git a/src/tint/sem/i32_type.cc b/src/tint/sem/i32_type.cc
index 33e12a0..dfe5dc3 100644
--- a/src/tint/sem/i32_type.cc
+++ b/src/tint/sem/i32_type.cc
@@ -27,6 +27,14 @@
 
 I32::~I32() = default;
 
+size_t I32::Hash() const {
+  return TypeInfo::Of<I32>().full_hashcode;
+}
+
+bool I32::Equals(const Type& other) const {
+  return other.Is<I32>();
+}
+
 std::string I32::type_name() const {
   return "__i32";
 }
diff --git a/src/tint/sem/i32_type.h b/src/tint/sem/i32_type.h
index 2128573..3700db0 100644
--- a/src/tint/sem/i32_type.h
+++ b/src/tint/sem/i32_type.h
@@ -31,6 +31,13 @@
   I32(I32&&);
   ~I32() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the name for this type
   std::string type_name() const override;
 
diff --git a/src/tint/sem/matrix_type.cc b/src/tint/sem/matrix_type.cc
index 8be69c3..0a28138 100644
--- a/src/tint/sem/matrix_type.cc
+++ b/src/tint/sem/matrix_type.cc
@@ -16,6 +16,7 @@
 
 #include "src/tint/program_builder.h"
 #include "src/tint/sem/vector_type.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::Matrix);
 
@@ -37,6 +38,19 @@
 
 Matrix::~Matrix() = default;
 
+size_t Matrix::Hash() const {
+  return utils::Hash(TypeInfo::Of<Vector>().full_hashcode, rows_, columns_,
+                     column_type_);
+}
+
+bool Matrix::Equals(const Type& other) const {
+  if (auto* v = other.As<Matrix>()) {
+    return v->rows_ == rows_ && v->columns_ == columns_ &&
+           v->column_type_ == column_type_;
+  }
+  return false;
+}
+
 std::string Matrix::type_name() const {
   return "__mat_" + std::to_string(rows_) + "_" + std::to_string(columns_) +
          subtype_->type_name();
diff --git a/src/tint/sem/matrix_type.h b/src/tint/sem/matrix_type.h
index bc5eb2a..b7adf68 100644
--- a/src/tint/sem/matrix_type.h
+++ b/src/tint/sem/matrix_type.h
@@ -36,6 +36,13 @@
   Matrix(Matrix&&);
   ~Matrix() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the type of the matrix
   const Type* type() const { return subtype_; }
   /// @returns the number of rows in the matrix
diff --git a/src/tint/sem/multisampled_texture_type.cc b/src/tint/sem/multisampled_texture_type.cc
index 717e5ba..4fe507f 100644
--- a/src/tint/sem/multisampled_texture_type.cc
+++ b/src/tint/sem/multisampled_texture_type.cc
@@ -15,6 +15,7 @@
 #include "src/tint/sem/multisampled_texture_type.h"
 
 #include "src/tint/program_builder.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::MultisampledTexture);
 
@@ -31,6 +32,18 @@
 
 MultisampledTexture::~MultisampledTexture() = default;
 
+size_t MultisampledTexture::Hash() const {
+  return utils::Hash(TypeInfo::Of<MultisampledTexture>().full_hashcode, dim(),
+                     type_);
+}
+
+bool MultisampledTexture::Equals(const sem::Type& other) const {
+  if (auto* o = other.As<MultisampledTexture>()) {
+    return o->dim() == dim() && o->type_ == type_;
+  }
+  return false;
+}
+
 std::string MultisampledTexture::type_name() const {
   std::ostringstream out;
   out << "__multisampled_texture_" << dim() << type_->type_name();
diff --git a/src/tint/sem/multisampled_texture_type.h b/src/tint/sem/multisampled_texture_type.h
index cb0f74b..f45671b 100644
--- a/src/tint/sem/multisampled_texture_type.h
+++ b/src/tint/sem/multisampled_texture_type.h
@@ -33,6 +33,13 @@
   MultisampledTexture(MultisampledTexture&&);
   ~MultisampledTexture() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the subtype of the sampled texture
   const Type* type() const { return type_; }
 
diff --git a/src/tint/sem/pointer_type.cc b/src/tint/sem/pointer_type.cc
index a548742..e73fc2e 100644
--- a/src/tint/sem/pointer_type.cc
+++ b/src/tint/sem/pointer_type.cc
@@ -16,6 +16,7 @@
 
 #include "src/tint/program_builder.h"
 #include "src/tint/sem/reference_type.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::Pointer);
 
@@ -30,6 +31,19 @@
   TINT_ASSERT(Semantic, access != ast::Access::kUndefined);
 }
 
+size_t Pointer::Hash() const {
+  return utils::Hash(TypeInfo::Of<Pointer>().full_hashcode, storage_class_,
+                     subtype_, access_);
+}
+
+bool Pointer::Equals(const sem::Type& other) const {
+  if (auto* o = other.As<Pointer>()) {
+    return o->storage_class_ == storage_class_ && o->subtype_ == subtype_ &&
+           o->access_ == access_;
+  }
+  return false;
+}
+
 std::string Pointer::type_name() const {
   std::ostringstream out;
   out << "__ptr_" << storage_class_ << subtype_->type_name() << "__" << access_;
diff --git a/src/tint/sem/pointer_type.h b/src/tint/sem/pointer_type.h
index 7bc8001..d940fce 100644
--- a/src/tint/sem/pointer_type.h
+++ b/src/tint/sem/pointer_type.h
@@ -39,6 +39,13 @@
   Pointer(Pointer&&);
   ~Pointer() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the pointee type
   const Type* StoreType() const { return subtype_; }
 
diff --git a/src/tint/sem/reference_type.cc b/src/tint/sem/reference_type.cc
index daa0984..1ad5487 100644
--- a/src/tint/sem/reference_type.cc
+++ b/src/tint/sem/reference_type.cc
@@ -15,6 +15,7 @@
 #include "src/tint/sem/reference_type.h"
 
 #include "src/tint/program_builder.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::Reference);
 
@@ -29,6 +30,19 @@
   TINT_ASSERT(Semantic, access != ast::Access::kUndefined);
 }
 
+size_t Reference::Hash() const {
+  return utils::Hash(TypeInfo::Of<Reference>().full_hashcode, storage_class_,
+                     subtype_, access_);
+}
+
+bool Reference::Equals(const sem::Type& other) const {
+  if (auto* o = other.As<Reference>()) {
+    return o->storage_class_ == storage_class_ && o->subtype_ == subtype_ &&
+           o->access_ == access_;
+  }
+  return false;
+}
+
 std::string Reference::type_name() const {
   std::ostringstream out;
   out << "__ref_" << storage_class_ << subtype_->type_name() << "__" << access_;
diff --git a/src/tint/sem/reference_type.h b/src/tint/sem/reference_type.h
index 67bc145..3af28d5 100644
--- a/src/tint/sem/reference_type.h
+++ b/src/tint/sem/reference_type.h
@@ -39,6 +39,13 @@
   Reference(Reference&&);
   ~Reference() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the pointee type
   const Type* StoreType() const { return subtype_; }
 
diff --git a/src/tint/sem/sampled_texture_type.cc b/src/tint/sem/sampled_texture_type.cc
index d14742e..c5785b9 100644
--- a/src/tint/sem/sampled_texture_type.cc
+++ b/src/tint/sem/sampled_texture_type.cc
@@ -15,6 +15,7 @@
 #include "src/tint/sem/sampled_texture_type.h"
 
 #include "src/tint/program_builder.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::SampledTexture);
 
@@ -30,6 +31,18 @@
 
 SampledTexture::~SampledTexture() = default;
 
+size_t SampledTexture::Hash() const {
+  return utils::Hash(TypeInfo::Of<SampledTexture>().full_hashcode, dim(),
+                     type_);
+}
+
+bool SampledTexture::Equals(const sem::Type& other) const {
+  if (auto* o = other.As<SampledTexture>()) {
+    return o->dim() == dim() && o->type_ == type_;
+  }
+  return false;
+}
+
 std::string SampledTexture::type_name() const {
   std::ostringstream out;
   out << "__sampled_texture_" << dim() << type_->type_name();
diff --git a/src/tint/sem/sampled_texture_type.h b/src/tint/sem/sampled_texture_type.h
index 34f3c6d..0f13c80 100644
--- a/src/tint/sem/sampled_texture_type.h
+++ b/src/tint/sem/sampled_texture_type.h
@@ -33,6 +33,13 @@
   SampledTexture(SampledTexture&&);
   ~SampledTexture() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the subtype of the sampled texture
   Type* type() const { return const_cast<Type*>(type_); }
 
diff --git a/src/tint/sem/sampler_type.cc b/src/tint/sem/sampler_type.cc
index 20c3a8a..f45e59d 100644
--- a/src/tint/sem/sampler_type.cc
+++ b/src/tint/sem/sampler_type.cc
@@ -15,6 +15,7 @@
 #include "src/tint/sem/sampler_type.h"
 
 #include "src/tint/program_builder.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::Sampler);
 
@@ -27,6 +28,17 @@
 
 Sampler::~Sampler() = default;
 
+size_t Sampler::Hash() const {
+  return utils::Hash(TypeInfo::Of<Sampler>().full_hashcode, kind_);
+}
+
+bool Sampler::Equals(const sem::Type& other) const {
+  if (auto* o = other.As<Sampler>()) {
+    return o->kind_ == kind_;
+  }
+  return false;
+}
+
 std::string Sampler::type_name() const {
   return std::string("__sampler_") +
          (kind_ == ast::SamplerKind::kSampler ? "sampler" : "comparison");
diff --git a/src/tint/sem/sampler_type.h b/src/tint/sem/sampler_type.h
index ca06991..e98b712 100644
--- a/src/tint/sem/sampler_type.h
+++ b/src/tint/sem/sampler_type.h
@@ -33,6 +33,13 @@
   Sampler(Sampler&&);
   ~Sampler() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the sampler type
   ast::SamplerKind kind() const { return kind_; }
 
diff --git a/src/tint/sem/sem_array_test.cc b/src/tint/sem/sem_array_test.cc
index 8d16c13..414f198 100644
--- a/src/tint/sem/sem_array_test.cc
+++ b/src/tint/sem/sem_array_test.cc
@@ -23,7 +23,7 @@
 
 TEST_F(ArrayTest, CreateSizedArray) {
   U32 u32;
-  auto* arr = create<Array>(&u32, 2, 4, 8, 32, 16);
+  auto* arr = create<Array>(&u32, 2u, 4u, 8u, 32u, 16u);
   EXPECT_EQ(arr->ElemType(), &u32);
   EXPECT_EQ(arr->Count(), 2u);
   EXPECT_EQ(arr->Align(), 4u);
@@ -36,7 +36,7 @@
 
 TEST_F(ArrayTest, CreateRuntimeArray) {
   U32 u32;
-  auto* arr = create<Array>(&u32, 0, 4, 8, 32, 32);
+  auto* arr = create<Array>(&u32, 0u, 4u, 8u, 32u, 32u);
   EXPECT_EQ(arr->ElemType(), &u32);
   EXPECT_EQ(arr->Count(), 0u);
   EXPECT_EQ(arr->Align(), 4u);
@@ -49,33 +49,33 @@
 
 TEST_F(ArrayTest, TypeName) {
   I32 i32;
-  auto* arr = create<Array>(&i32, 2, 0, 4, 4, 4);
+  auto* arr = create<Array>(&i32, 2u, 0u, 4u, 4u, 4u);
   EXPECT_EQ(arr->type_name(), "__array__i32_count_2_align_0_size_4_stride_4");
 }
 
 TEST_F(ArrayTest, FriendlyNameRuntimeSized) {
-  auto* arr = create<Array>(create<I32>(), 0, 0, 4, 4, 4);
+  auto* arr = create<Array>(create<I32>(), 0u, 0u, 4u, 4u, 4u);
   EXPECT_EQ(arr->FriendlyName(Symbols()), "array<i32>");
 }
 
 TEST_F(ArrayTest, FriendlyNameStaticSized) {
-  auto* arr = create<Array>(create<I32>(), 5, 4, 20, 4, 4);
+  auto* arr = create<Array>(create<I32>(), 5u, 4u, 20u, 4u, 4u);
   EXPECT_EQ(arr->FriendlyName(Symbols()), "array<i32, 5>");
 }
 
 TEST_F(ArrayTest, FriendlyNameRuntimeSizedNonImplicitStride) {
-  auto* arr = create<Array>(create<I32>(), 0, 0, 4, 8, 4);
+  auto* arr = create<Array>(create<I32>(), 0u, 0u, 4u, 8u, 4u);
   EXPECT_EQ(arr->FriendlyName(Symbols()), "@stride(8) array<i32>");
 }
 
 TEST_F(ArrayTest, FriendlyNameStaticSizedNonImplicitStride) {
-  auto* arr = create<Array>(create<I32>(), 5, 4, 20, 8, 4);
+  auto* arr = create<Array>(create<I32>(), 5u, 4u, 20u, 8u, 4u);
   EXPECT_EQ(arr->FriendlyName(Symbols()), "@stride(8) array<i32, 5>");
 }
 
 TEST_F(ArrayTest, TypeName_RuntimeArray) {
   I32 i32;
-  auto* arr = create<Array>(&i32, 2, 4, 8, 16, 16);
+  auto* arr = create<Array>(&i32, 2u, 4u, 8u, 16u, 16u);
   EXPECT_EQ(arr->type_name(), "__array__i32_count_2_align_4_size_8_stride_16");
 }
 
diff --git a/src/tint/sem/sem_struct_test.cc b/src/tint/sem/sem_struct_test.cc
index 9134aab..48428cf 100644
--- a/src/tint/sem/sem_struct_test.cc
+++ b/src/tint/sem/sem_struct_test.cc
@@ -28,8 +28,8 @@
       create<ast::Struct>(name, ast::StructMemberList{}, ast::AttributeList{});
   auto* ptr = impl;
   auto* s =
-      create<sem::Struct>(impl, impl->name, StructMemberList{}, 4 /* align */,
-                          8 /* size */, 16 /* size_no_padding */);
+      create<sem::Struct>(impl, impl->name, StructMemberList{}, 4u /* align */,
+                          8u /* size */, 16u /* size_no_padding */);
   EXPECT_EQ(s->Declaration(), ptr);
   EXPECT_EQ(s->Align(), 4u);
   EXPECT_EQ(s->Size(), 8u);
@@ -41,8 +41,8 @@
   auto* impl =
       create<ast::Struct>(name, ast::StructMemberList{}, ast::AttributeList{});
   auto* s =
-      create<sem::Struct>(impl, impl->name, StructMemberList{}, 4 /* align */,
-                          4 /* size */, 4 /* size_no_padding */);
+      create<sem::Struct>(impl, impl->name, StructMemberList{}, 4u /* align */,
+                          4u /* size */, 4u /* size_no_padding */);
   EXPECT_EQ(s->type_name(), "__struct_$1");
 }
 
@@ -51,8 +51,8 @@
   auto* impl =
       create<ast::Struct>(name, ast::StructMemberList{}, ast::AttributeList{});
   auto* s =
-      create<sem::Struct>(impl, impl->name, StructMemberList{}, 4 /* align */,
-                          4 /* size */, 4 /* size_no_padding */);
+      create<sem::Struct>(impl, impl->name, StructMemberList{}, 4u /* align */,
+                          4u /* size */, 4u /* size_no_padding */);
   EXPECT_EQ(s->FriendlyName(Symbols()), "my_struct");
 }
 
diff --git a/src/tint/sem/storage_texture_type.cc b/src/tint/sem/storage_texture_type.cc
index dbff8e8..0a8ddd4 100644
--- a/src/tint/sem/storage_texture_type.cc
+++ b/src/tint/sem/storage_texture_type.cc
@@ -15,6 +15,7 @@
 #include "src/tint/sem/storage_texture_type.h"
 
 #include "src/tint/program_builder.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::StorageTexture);
 
@@ -31,6 +32,19 @@
 
 StorageTexture::~StorageTexture() = default;
 
+size_t StorageTexture::Hash() const {
+  return utils::Hash(TypeInfo::Of<StorageTexture>().full_hashcode, dim(),
+                     texel_format_, access_);
+}
+
+bool StorageTexture::Equals(const sem::Type& other) const {
+  if (auto* o = other.As<StorageTexture>()) {
+    return o->dim() == dim() && o->texel_format_ == texel_format_ &&
+           o->access_ == access_;
+  }
+  return false;
+}
+
 std::string StorageTexture::type_name() const {
   std::ostringstream out;
   out << "__storage_texture_" << dim() << "_" << texel_format_ << "_"
diff --git a/src/tint/sem/storage_texture_type.h b/src/tint/sem/storage_texture_type.h
index 23c904d..4a9c767 100644
--- a/src/tint/sem/storage_texture_type.h
+++ b/src/tint/sem/storage_texture_type.h
@@ -43,6 +43,13 @@
   StorageTexture(StorageTexture&&);
   ~StorageTexture() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the storage subtype
   Type* type() const { return subtype_; }
 
diff --git a/src/tint/sem/struct.cc b/src/tint/sem/struct.cc
index 6168a9d..adec60d 100644
--- a/src/tint/sem/struct.cc
+++ b/src/tint/sem/struct.cc
@@ -21,6 +21,7 @@
 
 #include "src/tint/ast/struct_member.h"
 #include "src/tint/symbol_table.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::Struct);
 TINT_INSTANTIATE_TYPEINFO(tint::sem::StructMember);
@@ -51,6 +52,17 @@
 
 Struct::~Struct() = default;
 
+size_t Struct::Hash() const {
+  return utils::Hash(TypeInfo::Of<Struct>().full_hashcode, name_);
+}
+
+bool Struct::Equals(const sem::Type& other) const {
+  if (auto* o = other.As<Struct>()) {
+    return o->name_ == name_;
+  }
+  return false;
+}
+
 const StructMember* Struct::FindMember(Symbol name) const {
   for (auto* member : members_) {
     if (member->Declaration()->symbol == name) {
diff --git a/src/tint/sem/struct.h b/src/tint/sem/struct.h
index acb146e..abc1782 100644
--- a/src/tint/sem/struct.h
+++ b/src/tint/sem/struct.h
@@ -74,6 +74,13 @@
   /// Destructor
   ~Struct() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the struct
   const ast::Struct* Declaration() const { return declaration_; }
 
diff --git a/src/tint/sem/type.h b/src/tint/sem/type.h
index 97cc658..ed98c54 100644
--- a/src/tint/sem/type.h
+++ b/src/tint/sem/type.h
@@ -15,6 +15,7 @@
 #ifndef SRC_TINT_SEM_TYPE_H_
 #define SRC_TINT_SEM_TYPE_H_
 
+#include <functional>
 #include <string>
 
 #include "src/tint/sem/node.h"
@@ -37,6 +38,13 @@
   Type(Type&&);
   ~Type() override;
 
+  /// @returns a hash of the type.
+  virtual size_t Hash() const = 0;
+
+  /// @returns true if the this type is equal to the given type
+  virtual bool Equals(const Type&) const = 0;
+
+  /// [DEPRECATED]
   /// @returns the name for this type. The type name is unique over all types.
   virtual std::string type_name() const = 0;
 
@@ -115,4 +123,27 @@
 }  // namespace sem
 }  // namespace tint
 
+namespace std {
+
+/// std::hash specialization for tint::sem::Type
+template <>
+struct hash<tint::sem::Type> {
+  /// @param type the type to obtain a hash from
+  /// @returns the hash of the semantic type
+  size_t operator()(const tint::sem::Type& type) const { return type.Hash(); }
+};
+
+/// std::equal_to specialization for tint::sem::Type
+template <>
+struct equal_to<tint::sem::Type> {
+  /// @param a the first type to compare
+  /// @param b the second type to compare
+  /// @returns true if the two types are equal
+  bool operator()(const tint::sem::Type& a, const tint::sem::Type& b) const {
+    return a.Equals(b);
+  }
+};
+
+}  // namespace std
+
 #endif  // SRC_TINT_SEM_TYPE_H_
diff --git a/src/tint/sem/type_manager.h b/src/tint/sem/type_manager.h
index 7cdce66..50f5f2d 100644
--- a/src/tint/sem/type_manager.h
+++ b/src/tint/sem/type_manager.h
@@ -20,16 +20,15 @@
 #include <utility>
 
 #include "src/tint/sem/type.h"
-#include "src/tint/utils/block_allocator.h"
+#include "src/tint/utils/unique_allocator.h"
 
-namespace tint {
-namespace sem {
+namespace tint::sem {
 
 /// The type manager holds all the pointers to the known types.
-class Manager {
+class Manager : public utils::UniqueAllocator<Type> {
  public:
   /// Iterator is the type returned by begin() and end()
-  using Iterator = utils::BlockAllocator<sem::Type>::ConstIterator;
+  using Iterator = utils::BlockAllocator<Type>::ConstIterator;
 
   /// Constructor
   Manager();
@@ -45,24 +44,6 @@
   /// Destructor
   ~Manager();
 
-  /// Get the given type `T` from the type manager
-  /// @param args the arguments to pass to the type constructor
-  /// @return the pointer to the registered type
-  template <typename T, typename... ARGS>
-  T* Get(ARGS&&... args) {
-    // Note: We do not use std::forward here, as we may need to use the
-    // arguments again for the call to Create<T>() below.
-    auto name = T(args...).type_name();
-    auto it = by_name_.find(name);
-    if (it != by_name_.end()) {
-      return static_cast<T*>(it->second);
-    }
-
-    auto* type = types_.Create<T>(std::forward<ARGS>(args)...);
-    by_name_.emplace(name, type);
-    return type;
-  }
-
   /// Wrap returns a new Manager created with the types of `inner`.
   /// The Manager returned by Wrap is intended to temporarily extend the types
   /// of an existing immutable Manager.
@@ -74,27 +55,16 @@
   /// @return the Manager that wraps `inner`
   static Manager Wrap(const Manager& inner) {
     Manager out;
-    out.by_name_ = inner.by_name_;
+    out.items = inner.items;
     return out;
   }
 
-  /// Returns the type map
-  /// @returns the mapping from name string to type.
-  const std::unordered_map<std::string, sem::Type*>& types() const {
-    return by_name_;
-  }
-
   /// @returns an iterator to the beginning of the types
-  Iterator begin() const { return types_.Objects().begin(); }
+  Iterator begin() const { return allocator.Objects().begin(); }
   /// @returns an iterator to the end of the types
-  Iterator end() const { return types_.Objects().end(); }
-
- private:
-  std::unordered_map<std::string, sem::Type*> by_name_;
-  utils::BlockAllocator<sem::Type> types_;
+  Iterator end() const { return allocator.Objects().end(); }
 };
 
-}  // namespace sem
-}  // namespace tint
+}  // namespace tint::sem
 
 #endif  // SRC_TINT_SEM_TYPE_MANAGER_H_
diff --git a/src/tint/sem/u32_type.cc b/src/tint/sem/u32_type.cc
index 30cf2c9..f005267 100644
--- a/src/tint/sem/u32_type.cc
+++ b/src/tint/sem/u32_type.cc
@@ -27,6 +27,14 @@
 
 U32::U32(U32&&) = default;
 
+size_t U32::Hash() const {
+  return TypeInfo::Of<U32>().full_hashcode;
+}
+
+bool U32::Equals(const Type& other) const {
+  return other.Is<U32>();
+}
+
 std::string U32::type_name() const {
   return "__u32";
 }
diff --git a/src/tint/sem/u32_type.h b/src/tint/sem/u32_type.h
index af5706c..1b22764 100644
--- a/src/tint/sem/u32_type.h
+++ b/src/tint/sem/u32_type.h
@@ -31,6 +31,13 @@
   U32(U32&&);
   ~U32() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the name for th type
   std::string type_name() const override;
 
diff --git a/src/tint/sem/vector_type.cc b/src/tint/sem/vector_type.cc
index c0bc890..1a0c233 100644
--- a/src/tint/sem/vector_type.cc
+++ b/src/tint/sem/vector_type.cc
@@ -15,6 +15,7 @@
 #include "src/tint/sem/vector_type.h"
 
 #include "src/tint/program_builder.h"
+#include "src/tint/utils/hash.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::sem::Vector);
 
@@ -31,6 +32,17 @@
 
 Vector::~Vector() = default;
 
+size_t Vector::Hash() const {
+  return utils::Hash(TypeInfo::Of<Vector>().full_hashcode, width_, subtype_);
+}
+
+bool Vector::Equals(const Type& other) const {
+  if (auto* v = other.As<Vector>()) {
+    return v->width_ == width_ && v->subtype_ == subtype_;
+  }
+  return false;
+}
+
 std::string Vector::type_name() const {
   return "__vec_" + std::to_string(width_) + subtype_->type_name();
 }
diff --git a/src/tint/sem/vector_type.h b/src/tint/sem/vector_type.h
index cbce602..98686b0 100644
--- a/src/tint/sem/vector_type.h
+++ b/src/tint/sem/vector_type.h
@@ -33,6 +33,13 @@
   Vector(Vector&&);
   ~Vector() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the type of the vector elements
   const Type* type() const { return subtype_; }
 
diff --git a/src/tint/sem/vector_type_test.cc b/src/tint/sem/vector_type_test.cc
index c853b9b..9e8886d 100644
--- a/src/tint/sem/vector_type_test.cc
+++ b/src/tint/sem/vector_type_test.cc
@@ -30,13 +30,13 @@
 
 TEST_F(VectorTest, TypeName) {
   auto* i32 = create<I32>();
-  auto* v = create<Vector>(i32, 3);
+  auto* v = create<Vector>(i32, 3u);
   EXPECT_EQ(v->type_name(), "__vec_3__i32");
 }
 
 TEST_F(VectorTest, FriendlyName) {
   auto* f32 = create<F32>();
-  auto* v = create<Vector>(f32, 3);
+  auto* v = create<Vector>(f32, 3u);
   EXPECT_EQ(v->FriendlyName(Symbols()), "vec3<f32>");
 }
 
diff --git a/src/tint/sem/void_type.cc b/src/tint/sem/void_type.cc
index 4a7abed..0f14aaa 100644
--- a/src/tint/sem/void_type.cc
+++ b/src/tint/sem/void_type.cc
@@ -27,6 +27,14 @@
 
 Void::~Void() = default;
 
+size_t Void::Hash() const {
+  return TypeInfo::Of<Void>().full_hashcode;
+}
+
+bool Void::Equals(const Type& other) const {
+  return other.Is<Void>();
+}
+
 std::string Void::type_name() const {
   return "__void";
 }
diff --git a/src/tint/sem/void_type.h b/src/tint/sem/void_type.h
index 593029e..0e71346 100644
--- a/src/tint/sem/void_type.h
+++ b/src/tint/sem/void_type.h
@@ -31,6 +31,13 @@
   Void(Void&&);
   ~Void() override;
 
+  /// @returns a hash of the type.
+  size_t Hash() const override;
+
+  /// @param other the other type to compare against
+  /// @returns true if the this type is equal to the given type
+  bool Equals(const Type& other) const override;
+
   /// @returns the name for this type
   std::string type_name() const override;
 
diff --git a/src/tint/transform/transform_test.cc b/src/tint/transform/transform_test.cc
index 1114ac0..d166529 100644
--- a/src/tint/transform/transform_test.cc
+++ b/src/tint/transform/transform_test.cc
@@ -69,7 +69,7 @@
 
 TEST_F(CreateASTTypeForTest, Vector) {
   auto* vec = create([](ProgramBuilder& b) {
-    return b.create<sem::Vector>(b.create<sem::F32>(), 2);
+    return b.create<sem::Vector>(b.create<sem::F32>(), 2u);
   });
   ASSERT_TRUE(vec->Is<ast::Vector>());
   ASSERT_TRUE(vec->As<ast::Vector>()->type->Is<ast::F32>());
@@ -78,7 +78,7 @@
 
 TEST_F(CreateASTTypeForTest, ArrayImplicitStride) {
   auto* arr = create([](ProgramBuilder& b) {
-    return b.create<sem::Array>(b.create<sem::F32>(), 2, 4, 4, 32u, 32u);
+    return b.create<sem::Array>(b.create<sem::F32>(), 2u, 4u, 4u, 32u, 32u);
   });
   ASSERT_TRUE(arr->Is<ast::Array>());
   ASSERT_TRUE(arr->As<ast::Array>()->type->Is<ast::F32>());
@@ -91,7 +91,7 @@
 
 TEST_F(CreateASTTypeForTest, ArrayNonImplicitStride) {
   auto* arr = create([](ProgramBuilder& b) {
-    return b.create<sem::Array>(b.create<sem::F32>(), 2, 4, 4, 64u, 32u);
+    return b.create<sem::Array>(b.create<sem::F32>(), 2u, 4u, 4u, 64u, 32u);
   });
   ASSERT_TRUE(arr->Is<ast::Array>());
   ASSERT_TRUE(arr->As<ast::Array>()->type->Is<ast::F32>());
@@ -110,8 +110,8 @@
   auto* str = create([](ProgramBuilder& b) {
     auto* decl = b.Structure("S", {}, {});
     return b.create<sem::Struct>(decl, decl->name, sem::StructMemberList{},
-                                 4 /* align */, 4 /* size */,
-                                 4 /* size_no_padding */);
+                                 4u /* align */, 4u /* size */,
+                                 4u /* size_no_padding */);
   });
   ASSERT_TRUE(str->Is<ast::TypeName>());
   EXPECT_EQ(ast_type_builder.Symbols().NameFor(str->As<ast::TypeName>()->name),
diff --git a/src/tint/utils/unique_allocator.h b/src/tint/utils/unique_allocator.h
index 264ff4c..69242fb 100644
--- a/src/tint/utils/unique_allocator.h
+++ b/src/tint/utils/unique_allocator.h
@@ -34,37 +34,49 @@
   /// @return a pointer to an instance of `T` with the provided arguments.
   ///         If an existing instance of `T` has been constructed, then the same
   ///         pointer is returned.
-  template <typename... ARGS>
-  T* Get(ARGS&&... args) {
+  template <typename TYPE = T, typename... ARGS>
+  TYPE* Get(ARGS&&... args) {
     // Create a temporary T instance on the stack so that we can hash it, and
     // use it for equality lookup for the std::unordered_set. If the item is not
     // found in the set, then we create the persisted instance with the
     // allocator.
-    T key{args...};
+    TYPE key{args...};
     auto hash = HASH{}(key);
     auto it = items.find(Entry{hash, &key});
     if (it != items.end()) {
-      return it->ptr;
+      return static_cast<TYPE*>(it->ptr);
     }
-    auto* ptr = allocator.Create(std::forward<ARGS>(args)...);
+    auto* ptr = allocator.template Create<TYPE>(std::forward<ARGS>(args)...);
     items.emplace_hint(it, Entry{hash, ptr});
     return ptr;
   }
 
- private:
+ protected:
+  /// Entry is used as the entry to the unordered_set
   struct Entry {
+    /// The pre-calculated hash of the entry
     size_t hash;
+    /// Tge pointer to the unique object
     T* ptr;
   };
-  struct Hasher {
-    size_t operator()(Entry e) const { return e.hash; }
-  };
+  /// Comparator is the hashing and equality function used by the unordered_set
   struct Comparator {
+    /// Hashing function
+    /// @param e the entry
+    /// @returns the hash of the entry
+    size_t operator()(Entry e) const { return e.hash; }
+
+    /// Equality function
+    /// @param a the first entry to compare
+    /// @param b the second entry to compare
+    /// @returns true if the two entries are equal
     bool operator()(Entry a, Entry b) const { return EQUAL{}(*a.ptr, *b.ptr); }
   };
 
+  /// The block allocator used to allocate the unique objects
   BlockAllocator<T> allocator;
-  std::unordered_set<Entry, Hasher, Comparator> items;
+  /// The unordered_set of unique item entries
+  std::unordered_set<Entry, Comparator, Comparator> items;
 };
 
 }  // namespace tint::utils
diff --git a/src/tint/writer/glsl/generator_impl_type_test.cc b/src/tint/writer/glsl/generator_impl_type_test.cc
index 64547e3..bc08fd2 100644
--- a/src/tint/writer/glsl/generator_impl_type_test.cc
+++ b/src/tint/writer/glsl/generator_impl_type_test.cc
@@ -122,8 +122,8 @@
 
 TEST_F(GlslGeneratorImplTest_Type, EmitType_Matrix) {
   auto* f32 = create<sem::F32>();
-  auto* vec3 = create<sem::Vector>(f32, 3);
-  auto* mat2x3 = create<sem::Matrix>(vec3, 2);
+  auto* vec3 = create<sem::Vector>(f32, 3u);
+  auto* mat2x3 = create<sem::Matrix>(vec3, 2u);
 
   GeneratorImpl& gen = Build();
 
@@ -224,7 +224,7 @@
 
 TEST_F(GlslGeneratorImplTest_Type, EmitType_Vector) {
   auto* f32 = create<sem::F32>();
-  auto* vec3 = create<sem::Vector>(f32, 3);
+  auto* vec3 = create<sem::Vector>(f32, 3u);
 
   GeneratorImpl& gen = Build();
 
diff --git a/src/tint/writer/hlsl/generator_impl_type_test.cc b/src/tint/writer/hlsl/generator_impl_type_test.cc
index 4e4ae32..f79164d 100644
--- a/src/tint/writer/hlsl/generator_impl_type_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_type_test.cc
@@ -122,8 +122,8 @@
 
 TEST_F(HlslGeneratorImplTest_Type, EmitType_Matrix) {
   auto* f32 = create<sem::F32>();
-  auto* vec3 = create<sem::Vector>(f32, 3);
-  auto* mat2x3 = create<sem::Matrix>(vec3, 2);
+  auto* vec3 = create<sem::Vector>(f32, 3u);
+  auto* mat2x3 = create<sem::Matrix>(vec3, 2u);
 
   GeneratorImpl& gen = Build();
 
@@ -241,7 +241,7 @@
 
 TEST_F(HlslGeneratorImplTest_Type, EmitType_Vector) {
   auto* f32 = create<sem::F32>();
-  auto* vec3 = create<sem::Vector>(f32, 3);
+  auto* vec3 = create<sem::Vector>(f32, 3u);
 
   GeneratorImpl& gen = Build();
 
diff --git a/src/tint/writer/msl/generator_impl_type_test.cc b/src/tint/writer/msl/generator_impl_type_test.cc
index d5de70e..57a9e59 100644
--- a/src/tint/writer/msl/generator_impl_type_test.cc
+++ b/src/tint/writer/msl/generator_impl_type_test.cc
@@ -177,8 +177,8 @@
 
 TEST_F(MslGeneratorImplTest, EmitType_Matrix) {
   auto* f32 = create<sem::F32>();
-  auto* vec3 = create<sem::Vector>(f32, 3);
-  auto* mat2x3 = create<sem::Matrix>(vec3, 2);
+  auto* vec3 = create<sem::Vector>(f32, 3u);
+  auto* mat2x3 = create<sem::Matrix>(vec3, 2u);
 
   GeneratorImpl& gen = Build();
 
@@ -708,7 +708,7 @@
 
 TEST_F(MslGeneratorImplTest, EmitType_Vector) {
   auto* f32 = create<sem::F32>();
-  auto* vec3 = create<sem::Vector>(f32, 3);
+  auto* vec3 = create<sem::Vector>(f32, 3u);
 
   GeneratorImpl& gen = Build();
 
diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc
index bec89bd..2884b3e 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -516,7 +516,7 @@
       has_overridable_workgroup_size_ = true;
 
       auto* vec3_u32 =
-          builder_.create<sem::Vector>(builder_.create<sem::U32>(), 3);
+          builder_.create<sem::Vector>(builder_.create<sem::U32>(), 3u);
       uint32_t vec3_u32_type_id = GenerateTypeIfNeeded(vec3_u32);
       if (vec3_u32_type_id == 0) {
         return 0;
@@ -2712,7 +2712,7 @@
     if (texture_type
             ->IsAnyOf<sem::DepthTexture, sem::DepthMultisampledTexture>()) {
       auto* f32 = builder_.create<sem::F32>();
-      auto* spirv_result_type = builder_.create<sem::Vector>(f32, 4);
+      auto* spirv_result_type = builder_.create<sem::Vector>(f32, 4u);
       auto spirv_result = result_op();
       post_emission = [=] {
         return push_function_inst(
@@ -3996,9 +3996,7 @@
       [&](const sem::StorageTexture* t) {
         return GenerateTypeIfNeeded(t->type());
       },
-      [&](Default)  {
-        return 0u;
-      });
+      [&](Default) { return 0u; });
   if (type_id == 0u) {
     return false;
   }
diff --git a/src/tint/writer/spirv/builder_type_test.cc b/src/tint/writer/spirv/builder_type_test.cc
index 7fecdc1..84d211c 100644
--- a/src/tint/writer/spirv/builder_type_test.cc
+++ b/src/tint/writer/spirv/builder_type_test.cc
@@ -208,8 +208,8 @@
 
 TEST_F(BuilderTest_Type, GenerateMatrix) {
   auto* f32 = create<sem::F32>();
-  auto* vec3 = create<sem::Vector>(f32, 3);
-  auto* mat2x3 = create<sem::Matrix>(vec3, 2);
+  auto* vec3 = create<sem::Vector>(f32, 3u);
+  auto* mat2x3 = create<sem::Matrix>(vec3, 2u);
 
   spirv::Builder& b = Build();
 
@@ -226,8 +226,8 @@
 
 TEST_F(BuilderTest_Type, ReturnsGeneratedMatrix) {
   auto* i32 = create<sem::I32>();
-  auto* col = create<sem::Vector>(i32, 4);
-  auto* mat = create<sem::Matrix>(col, 3);
+  auto* col = create<sem::Vector>(i32, 4u);
+  auto* mat = create<sem::Matrix>(col, 3u);
 
   spirv::Builder& b = Build();
 
@@ -469,7 +469,7 @@
 }
 
 TEST_F(BuilderTest_Type, GenerateVector) {
-  auto* vec = create<sem::Vector>(create<sem::F32>(), 3);
+  auto* vec = create<sem::Vector>(create<sem::F32>(), 3u);
 
   spirv::Builder& b = Build();
 
@@ -485,7 +485,7 @@
 
 TEST_F(BuilderTest_Type, ReturnsGeneratedVector) {
   auto* i32 = create<sem::I32>();
-  auto* vec = create<sem::Vector>(i32, 3);
+  auto* vec = create<sem::Vector>(i32, 3u);
 
   spirv::Builder& b = Build();