diff --git a/docs/origin-trial-changes.md b/docs/origin-trial-changes.md
index d5d69f1..3437a84 100644
--- a/docs/origin-trial-changes.md
+++ b/docs/origin-trial-changes.md
@@ -12,6 +12,7 @@
 
 ### New Features
 
+* Shadowing is now fully supported. [tint:819](https://crbug.com/tint/819)
 * The `dot()` builtin now supports integer vector types.
 * Identifiers can now start with a single leading underscore.  [tint:1292](https://crbug.com/tint/1292)
 
diff --git a/src/BUILD.gn b/src/BUILD.gn
index b225b19..2fd70f5 100644
--- a/src/BUILD.gn
+++ b/src/BUILD.gn
@@ -474,6 +474,8 @@
     "transform/single_entry_point.h",
     "transform/transform.cc",
     "transform/transform.h",
+    "transform/unshadow.cc",
+    "transform/unshadow.h",
     "transform/vectorize_scalar_matrix_constructors.cc",
     "transform/vectorize_scalar_matrix_constructors.h",
     "transform/vertex_pulling.cc",
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 8d09834..0e8c98e 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -338,6 +338,8 @@
   transform/single_entry_point.h
   transform/transform.cc
   transform/transform.h
+  transform/unshadow.cc
+  transform/unshadow.h
   transform/vectorize_scalar_matrix_constructors.cc
   transform/vectorize_scalar_matrix_constructors.h
   transform/vertex_pulling.cc
@@ -969,7 +971,8 @@
       transform/simplify_pointers_test.cc
       transform/single_entry_point_test.cc
       transform/test_helper.h
-      transform/vectorize_scalar_matrix_constructors.cc
+      transform/unshadow_test.cc
+      transform/vectorize_scalar_matrix_constructors_test.cc
       transform/vertex_pulling_test.cc
       transform/wrap_arrays_in_structs_test.cc
       transform/zero_init_workgroup_memory_test.cc
diff --git a/src/reader/spirv/parser.cc b/src/reader/spirv/parser.cc
index c5ffebb..6500966 100644
--- a/src/reader/spirv/parser.cc
+++ b/src/reader/spirv/parser.cc
@@ -20,6 +20,7 @@
 #include "src/transform/decompose_strided_matrix.h"
 #include "src/transform/manager.h"
 #include "src/transform/simplify_pointers.h"
+#include "src/transform/unshadow.h"
 
 namespace tint {
 namespace reader {
@@ -52,6 +53,7 @@
   // attribute then we need to decompose these into an array of vectors
   if (transform::DecomposeStridedMatrix::ShouldRun(&program)) {
     transform::Manager manager;
+    manager.Add<transform::Unshadow>();
     manager.Add<transform::SimplifyPointers>();
     manager.Add<transform::DecomposeStridedMatrix>();
     return manager.Run(&program).program;
diff --git a/src/resolver/dependency_graph.cc b/src/resolver/dependency_graph.cc
index 2de794f..4cfb45d 100644
--- a/src/resolver/dependency_graph.cc
+++ b/src/resolver/dependency_graph.cc
@@ -177,6 +177,9 @@
     TINT_DEFER(scope_stack_.Pop());
 
     for (auto* param : func->params) {
+      if (auto* shadows = scope_stack_.Get(param->symbol)) {
+        graph_.shadows.emplace(param, shadows);
+      }
       Declare(param->symbol, param);
       TraverseType(param->type);
     }
@@ -255,6 +258,9 @@
       return;
     }
     if (auto* v = stmt->As<ast::VariableDeclStatement>()) {
+      if (auto* shadows = scope_stack_.Get(v->variable->symbol)) {
+        graph_.shadows.emplace(v->variable, shadows);
+      }
       TraverseType(v->variable->type);
       TraverseExpression(v->variable->constructor);
       Declare(v->variable->symbol, v->variable);
diff --git a/src/resolver/dependency_graph.h b/src/resolver/dependency_graph.h
index b97e3fd..6235ad7 100644
--- a/src/resolver/dependency_graph.h
+++ b/src/resolver/dependency_graph.h
@@ -55,6 +55,12 @@
   /// Map of ast::IdentifierExpression or ast::TypeName to a type, function, or
   /// variable that declares the symbol.
   std::unordered_map<const ast::Node*, const ast::Node*> resolved_symbols;
+
+  /// Map of ast::Variable to a type, function, or variable that is shadowed by
+  /// the variable key. A declaration (X) shadows another (Y) if X and Y use
+  /// the same symbol, and X is declared in a sub-scope of the scope that
+  /// declares Y.
+  std::unordered_map<const ast::Variable*, const ast::Node*> shadows;
 };
 
 }  // namespace resolver
diff --git a/src/resolver/dependency_graph_test.cc b/src/resolver/dependency_graph_test.cc
index d256197..e222b43 100644
--- a/src/resolver/dependency_graph_test.cc
+++ b/src/resolver/dependency_graph_test.cc
@@ -88,6 +88,17 @@
     SymbolDeclKind::NestedLocalLet,
 };
 
+static constexpr SymbolDeclKind kGlobalDeclKinds[] = {
+    SymbolDeclKind::GlobalVar, SymbolDeclKind::GlobalLet, SymbolDeclKind::Alias,
+    SymbolDeclKind::Struct,    SymbolDeclKind::Function,
+};
+
+static constexpr SymbolDeclKind kLocalDeclKinds[] = {
+    SymbolDeclKind::Parameter,      SymbolDeclKind::LocalVar,
+    SymbolDeclKind::LocalLet,       SymbolDeclKind::NestedLocalVar,
+    SymbolDeclKind::NestedLocalLet,
+};
+
 static constexpr SymbolDeclKind kGlobalValueDeclKinds[] = {
     SymbolDeclKind::GlobalVar,
     SymbolDeclKind::GlobalLet,
@@ -1172,6 +1183,53 @@
 }  // namespace resolved_symbols
 
 ////////////////////////////////////////////////////////////////////////////////
+// Shadowing tests
+////////////////////////////////////////////////////////////////////////////////
+namespace shadowing {
+
+using ResolverDependencyShadowTest = ResolverDependencyGraphTestWithParam<
+    std::tuple<SymbolDeclKind, SymbolDeclKind>>;
+
+TEST_P(ResolverDependencyShadowTest, Test) {
+  const Symbol symbol = Sym("SYMBOL");
+  const auto outer_kind = std::get<0>(GetParam());
+  const auto inner_kind = std::get<1>(GetParam());
+
+  // Build a symbol declaration and a use of that symbol
+  SymbolTestHelper helper(this);
+  auto* outer = helper.Add(outer_kind, symbol, Source{{12, 34}});
+  helper.Add(inner_kind, symbol, Source{{56, 78}});
+  auto* inner_var = helper.nested_statements.size()
+                        ? helper.nested_statements[0]
+                              ->As<ast::VariableDeclStatement>()
+                              ->variable
+                        : helper.statements.size()
+                              ? helper.statements[0]
+                                    ->As<ast::VariableDeclStatement>()
+                                    ->variable
+                              : helper.parameters[0];
+  helper.Build();
+
+  EXPECT_EQ(Build().shadows[inner_var], outer);
+}
+
+INSTANTIATE_TEST_SUITE_P(LocalShadowGlobal,
+                         ResolverDependencyShadowTest,
+                         testing::Combine(testing::ValuesIn(kGlobalDeclKinds),
+                                          testing::ValuesIn(kLocalDeclKinds)));
+
+INSTANTIATE_TEST_SUITE_P(
+    NestedLocalShadowLocal,
+    ResolverDependencyShadowTest,
+    testing::Combine(testing::Values(SymbolDeclKind::Parameter,
+                                     SymbolDeclKind::LocalVar,
+                                     SymbolDeclKind::LocalLet),
+                     testing::Values(SymbolDeclKind::NestedLocalVar,
+                                     SymbolDeclKind::NestedLocalLet)));
+
+}  // namespace shadowing
+
+////////////////////////////////////////////////////////////////////////////////
 // AST traversal tests
 ////////////////////////////////////////////////////////////////////////////////
 namespace ast_traversal {
diff --git a/src/resolver/function_validation_test.cc b/src/resolver/function_validation_test.cc
index 09dab00..af3b097 100644
--- a/src/resolver/function_validation_test.cc
+++ b/src/resolver/function_validation_test.cc
@@ -26,25 +26,46 @@
 class ResolverFunctionValidationTest : public resolver::TestHelper,
                                        public testing::Test {};
 
-TEST_F(ResolverFunctionValidationTest, ParameterNamesMustBeUnique_pass) {
+TEST_F(ResolverFunctionValidationTest, DuplicateParameterName) {
   // fn func_a(common_name : f32) { }
   // fn func_b(common_name : f32) { }
   Func("func_a", {Param("common_name", ty.f32())}, ty.void_(), {});
   Func("func_b", {Param("common_name", ty.f32())}, ty.void_(), {});
 
-  EXPECT_TRUE(r()->Resolve());
-  EXPECT_EQ(r()->error(), "");
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
-TEST_F(ResolverFunctionValidationTest,
-       ParameterNamesMustBeUniqueShadowsGlobal_pass) {
+TEST_F(ResolverFunctionValidationTest, ParameterMayShadowGlobal) {
   // var<private> common_name : f32;
   // fn func(common_name : f32) { }
   Global("common_name", ty.f32(), ast::StorageClass::kPrivate);
   Func("func", {Param("common_name", ty.f32())}, ty.void_(), {});
 
-  EXPECT_TRUE(r()->Resolve());
-  EXPECT_EQ(r()->error(), "");
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverFunctionValidationTest, LocalConflictsWithParameter) {
+  // fn func(common_name : f32) {
+  //   let common_name = 1;
+  // }
+  Func("func", {Param(Source{{12, 34}}, "common_name", ty.f32())}, ty.void_(),
+       {Decl(Const(Source{{56, 78}}, "common_name", nullptr, Expr(1)))});
+
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(), R"(56:78 error: redeclaration of 'common_name'
+12:34 note: 'common_name' previously declared here)");
+}
+
+TEST_F(ResolverFunctionValidationTest, NestedLocalMayShadowParameter) {
+  // fn func(common_name : f32) {
+  //   {
+  //     let common_name = 1;
+  //   }
+  // }
+  Func("func", {Param(Source{{12, 34}}, "common_name", ty.f32())}, ty.void_(),
+       {Block(Decl(Const(Source{{56, 78}}, "common_name", nullptr, Expr(1))))});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest,
@@ -57,7 +78,7 @@
            Decl(var),
        });
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest, FunctionUsingSameVariableName_Pass) {
@@ -74,7 +95,7 @@
        },
        ast::DecorationList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest,
@@ -95,7 +116,7 @@
        },
        ast::DecorationList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest, UnreachableCode_return) {
@@ -189,7 +210,7 @@
   Func(Source{{12, 34}}, "func", ast::VariableList{}, ty.void_(),
        ast::StatementList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest,
@@ -213,7 +234,7 @@
            Return(),
        });
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest,
@@ -269,7 +290,7 @@
        },
        ast::DecorationList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest,
@@ -298,7 +319,7 @@
        },
        ast::DecorationList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest,
@@ -362,7 +383,7 @@
        },
        ast::DecorationList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest, FunctionVarInitWithParam) {
@@ -376,7 +397,7 @@
   Func("foo", ast::VariableList{bar}, ty.void_(), ast::StatementList{Decl(baz)},
        ast::DecorationList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest, FunctionConstInitWithParam) {
@@ -390,7 +411,7 @@
   Func("foo", ast::VariableList{bar}, ty.void_(), ast::StatementList{Decl(baz)},
        ast::DecorationList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest, FunctionParamsConst) {
@@ -414,7 +435,7 @@
        {Stage(ast::PipelineStage::kCompute),
         WorkgroupSize(Expr("x"), Expr("y"), Expr(16u))});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_U32) {
@@ -425,7 +446,7 @@
        {Stage(ast::PipelineStage::kCompute),
         WorkgroupSize(Source{{12, 34}}, Expr(1u), Expr(2u), Expr(3u))});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest, WorkgroupSize_MismatchTypeU32) {
@@ -696,7 +717,7 @@
   auto* bar = Param(Source{{12, 34}}, "bar", ret_type);
   Func("f", ast::VariableList{bar}, ty.void_(), {});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest, ParametersAtLimit) {
@@ -706,7 +727,7 @@
   }
   Func(Source{{12, 34}}, "f", params, ty.void_(), {});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest, ParametersOverLimit) {
@@ -736,7 +757,7 @@
   Func("f", ast::VariableList{arg}, ty.void_(), {});
 
   if (param.should_pass) {
-    EXPECT_TRUE(r()->Resolve()) << r()->error();
+    ASSERT_TRUE(r()->Resolve()) << r()->error();
   } else {
     std::stringstream ss;
     ss << param.storage_class;
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc
index b075644..0e0abc9 100644
--- a/src/resolver/resolver.cc
+++ b/src/resolver/resolver.cc
@@ -74,7 +74,6 @@
 #include "src/sem/type_conversion.h"
 #include "src/sem/variable.h"
 #include "src/utils/defer.h"
-#include "src/utils/map.h"
 #include "src/utils/math.h"
 #include "src/utils/reverse.h"
 #include "src/utils/scoped_assignment.h"
@@ -142,6 +141,8 @@
 
   AllocateOverridableConstantIds();
 
+  SetShadows();
+
   if (!ValidatePipelineStages()) {
     return false;
   }
@@ -258,14 +259,8 @@
     if (ty->As<ast::ExternalTexture>()) {
       return builder_->create<sem::ExternalTexture>();
     }
-    if (auto* t = ty->As<ast::TypeName>()) {
-      auto it = named_type_info_.find(t->name);
-      if (it == named_type_info_.end()) {
-        AddError("unknown type '" + builder_->Symbols().NameFor(t->name) + "'",
-                 t->source);
-        return nullptr;
-      }
-      return it->second.sem;
+    if (auto* type = ResolvedSymbol<sem::Type>(ty)) {
+      return type;
     }
     TINT_UNREACHABLE(Resolver, diagnostics_)
         << "Unhandled ast::Type: " << ty->TypeInfo().name;
@@ -423,7 +418,7 @@
     }
     case VariableKind::kLocal: {
       auto* local = builder_->create<sem::LocalVariable>(
-          var, var_ty, storage_class, access,
+          var, var_ty, storage_class, access, current_statement_,
           (rhs && var->is_const) ? rhs->ConstantValue() : sem::Constant{});
       builder_->Sem().Add(var, local);
       return local;
@@ -496,17 +491,23 @@
   }
 }
 
-bool Resolver::GlobalVariable(const ast::Variable* var) {
-  if (!ValidateNoDuplicateDefinition(var->symbol, var->source,
-                                     /* check_global_scope_only */ true)) {
-    return false;
+void Resolver::SetShadows() {
+  for (auto it : dependencies_.shadows) {
+    auto* var = Sem(it.first);
+    if (auto* local = var->As<sem::LocalVariable>()) {
+      local->SetShadows(Sem(it.second));
+    }
+    if (auto* param = var->As<sem::Parameter>()) {
+      param->SetShadows(Sem(it.second));
+    }
   }
+}  // namespace resolver
 
+bool Resolver::GlobalVariable(const ast::Variable* var) {
   auto* sem = Variable(var, VariableKind::kGlobal);
   if (!sem) {
     return false;
   }
-  variable_stack_.Set(var->symbol, sem);
 
   auto storage_class = sem->StorageClass();
   if (!var->is_const && storage_class == ast::StorageClass::kNone) {
@@ -547,9 +548,6 @@
 }
 
 sem::Function* Resolver::Function(const ast::Function* decl) {
-  variable_stack_.Push();
-  TINT_DEFER(variable_stack_.Pop());
-
   uint32_t parameter_index = 0;
   std::unordered_map<Symbol, Source> parameter_names;
   std::vector<sem::Parameter*> parameters;
@@ -581,7 +579,6 @@
       return nullptr;
     }
 
-    variable_stack_.Set(param->symbol, var);
     parameters.emplace_back(var);
 
     auto* var_ty = const_cast<sem::Type*>(var->Type());
@@ -684,11 +681,6 @@
     return nullptr;
   }
 
-  // Register the function information _after_ processing the statements. This
-  // allows us to catch a function calling itself when determining the call
-  // information as this function doesn't exist until it's finished.
-  symbol_to_function_[decl->symbol] = func;
-
   // If this is an entry point, mark all transitively called functions as being
   // used by this entry point.
   if (decl->IsEntryPoint()) {
@@ -1241,22 +1233,28 @@
   auto* ident = expr->target.name;
   Mark(ident);
 
-  auto it = named_type_info_.find(ident->symbol);
-  if (it != named_type_info_.end()) {
-    // We have a type.
-    return type_ctor_or_conv(it->second.sem);
+  auto* resolved = ResolvedSymbol(ident);
+  if (auto* ty = As<sem::Type>(resolved)) {
+    return type_ctor_or_conv(ty);
   }
 
-  // Not a type, treat as a intrinsic / function call.
+  if (auto* fn = As<sem::Function>(resolved)) {
+    return FunctionCall(expr, fn, std::move(args));
+  }
+
   auto name = builder_->Symbols().NameFor(ident->symbol);
   auto intrinsic_type = sem::ParseIntrinsicType(name);
-  auto* call = (intrinsic_type != sem::IntrinsicType::kNone)
-                   ? IntrinsicCall(expr, intrinsic_type, std::move(args),
-                                   std::move(arg_tys))
-                   : FunctionCall(expr, std::move(args));
+  if (intrinsic_type != sem::IntrinsicType::kNone) {
+    return IntrinsicCall(expr, intrinsic_type, std::move(args),
+                         std::move(arg_tys));
+  }
 
-  current_function_->AddDirectCall(call);
-  return call;
+  TINT_ICE(Resolver, diagnostics_)
+      << expr->source << " unresolved CallExpression target:\n"
+      << "resolved: " << (resolved ? resolved->TypeInfo().name : "<null>")
+      << "\n"
+      << "name: " << builder_->Symbols().NameFor(ident->symbol);
+  return nullptr;
 }
 
 sem::Call* Resolver::IntrinsicCall(
@@ -1288,37 +1286,27 @@
     return nullptr;
   }
 
+  current_function_->AddDirectCall(call);
+
   return call;
 }
 
 sem::Call* Resolver::FunctionCall(
     const ast::CallExpression* expr,
+    sem::Function* target,
     const std::vector<const sem::Expression*> args) {
   auto sym = expr->target.name->symbol;
   auto name = builder_->Symbols().NameFor(sym);
 
-  auto target_it = symbol_to_function_.find(sym);
-  if (target_it == symbol_to_function_.end()) {
-    if (current_function_ && current_function_->Declaration()->symbol == sym) {
-      AddError("recursion is not permitted. '" + name +
-                   "' attempted to call itself.",
-               expr->source);
-    } else {
-      AddError("unable to find called function: " + name, expr->source);
-    }
-    return nullptr;
-  }
-  auto* target = target_it->second;
   auto* call = builder_->create<sem::Call>(expr, target, std::move(args),
                                            current_statement_, sem::Constant{});
 
   if (current_function_) {
-    target->AddCallSite(call);
-
     // Note: Requires called functions to be resolved first.
     // This is currently guaranteed as functions must be declared before
     // use.
     current_function_->AddTransitivelyCalledFunction(target);
+    current_function_->AddDirectCall(call);
     for (auto* transitive_call : target->TransitivelyCalledFunctions()) {
       current_function_->AddTransitivelyCalledFunction(transitive_call);
     }
@@ -1329,6 +1317,8 @@
     }
   }
 
+  target->AddCallSite(call);
+
   if (!ValidateFunctionCall(call)) {
     return nullptr;
   }
@@ -1476,7 +1466,8 @@
 
 sem::Expression* Resolver::Identifier(const ast::IdentifierExpression* expr) {
   auto symbol = expr->symbol;
-  if (auto* var = variable_stack_.Get(symbol)) {
+  auto* resolved = ResolvedSymbol(expr);
+  if (auto* var = As<sem::Variable>(resolved)) {
     auto* user =
         builder_->create<sem::VariableUser>(expr, current_statement_, var);
 
@@ -1526,18 +1517,26 @@
     return user;
   }
 
-  if (symbol_to_function_.count(symbol)) {
+  if (Is<sem::Function>(resolved)) {
     AddError("missing '(' for function call", expr->source.End());
     return nullptr;
   }
 
-  std::string name = builder_->Symbols().NameFor(symbol);
-  if (sem::ParseIntrinsicType(name) != sem::IntrinsicType::kNone) {
+  if (IsIntrinsic(symbol)) {
     AddError("missing '(' for intrinsic call", expr->source.End());
     return nullptr;
   }
 
-  AddError("unknown identifier: '" + name + "'", expr->source);
+  if (resolved->Is<sem::Type>()) {
+    AddError("missing '(' for type constructor or cast", expr->source.End());
+    return nullptr;
+  }
+
+  TINT_ICE(Resolver, diagnostics_)
+      << expr->source << " unresolved identifier:\n"
+      << "resolved: " << (resolved ? resolved->TypeInfo().name : "<null>")
+      << "\n"
+      << "name: " << builder_->Symbols().NameFor(symbol);
   return nullptr;
 }
 
@@ -1934,11 +1933,6 @@
 bool Resolver::VariableDeclStatement(const ast::VariableDeclStatement* stmt) {
   Mark(stmt->variable);
 
-  if (!ValidateNoDuplicateDefinition(stmt->variable->symbol,
-                                     stmt->variable->source)) {
-    return false;
-  }
-
   auto* var = Variable(stmt->variable, VariableKind::kLocal);
   if (!var) {
     return false;
@@ -1952,7 +1946,6 @@
     }
   }
 
-  variable_stack_.Set(stmt->variable->symbol, var);
   if (current_block_) {  // Not all statements are inside a block
     current_block_->AddDecl(stmt->variable);
   }
@@ -1978,12 +1971,6 @@
     return nullptr;
   }
 
-  named_type_info_.emplace(named_type->name, TypeDeclInfo{named_type, result});
-
-  if (!ValidateTypeDecl(named_type)) {
-    return nullptr;
-  }
-
   builder_->Sem().Add(named_type, result);
   return result;
 }
@@ -2081,7 +2068,7 @@
 
     if (auto* ident = count_expr->As<ast::IdentifierExpression>()) {
       // Make sure the identifier is a non-overridable module-scope constant.
-      auto* var = variable_stack_.Get(ident->symbol);
+      auto* var = ResolvedSymbol<sem::Variable>(ident);
       if (!var || !var->Is<sem::GlobalVariable>() ||
           !var->Declaration()->is_const) {
         AddError("array size identifier must be a module-scope constant",
@@ -2244,13 +2231,11 @@
   current_statement_ = stmt;
   current_compound_statement_ = stmt;
   current_block_ = stmt->As<sem::BlockStatement>();
-  variable_stack_.Push();
 
   TINT_DEFER({
     current_block_ = prev_current_block;
     current_compound_statement_ = prev_current_compound_statement;
     current_statement_ = prev_current_statement;
-    variable_stack_.Pop();
   });
 
   return callback();
@@ -2330,6 +2315,11 @@
   return false;
 }
 
+bool Resolver::IsIntrinsic(Symbol symbol) const {
+  std::string name = builder_->Symbols().NameFor(symbol);
+  return sem::ParseIntrinsicType(name) != sem::IntrinsicType::kNone;
+}
+
 ////////////////////////////////////////////////////////////////////////////////
 // Resolver::TypeConversionSig
 ////////////////////////////////////////////////////////////////////////////////
diff --git a/src/resolver/resolver.h b/src/resolver/resolver.h
index a55c38f..cca4440 100644
--- a/src/resolver/resolver.h
+++ b/src/resolver/resolver.h
@@ -32,6 +32,7 @@
 #include "src/sem/constant.h"
 #include "src/sem/function.h"
 #include "src/sem/struct.h"
+#include "src/utils/map.h"
 #include "src/utils/unique_vector.h"
 
 namespace tint {
@@ -176,6 +177,7 @@
   sem::Expression* Expression(const ast::Expression*);
   sem::Function* Function(const ast::Function*);
   sem::Call* FunctionCall(const ast::CallExpression*,
+                          sem::Function* target,
                           const std::vector<const sem::Expression*> args);
   sem::Expression* Identifier(const ast::IdentifierExpression*);
   sem::Call* IntrinsicCall(const ast::CallExpression*,
@@ -239,9 +241,6 @@
   bool ValidateMatrix(const sem::Matrix* ty, const Source& source);
   bool ValidateFunctionParameter(const ast::Function* func,
                                  const sem::Variable* var);
-  bool ValidateNoDuplicateDefinition(Symbol sym,
-                                     const Source& source,
-                                     bool check_global_scope_only = false);
   bool ValidateParameter(const ast::Function* func, const sem::Variable* var);
   bool ValidateReturn(const ast::ReturnStatement* ret);
   bool ValidateStatements(const ast::StatementList& stmts);
@@ -264,7 +263,6 @@
                                        const sem::Type* type);
   bool ValidateArrayConstructorOrCast(const ast::CallExpression* ctor,
                                       const sem::Array* arr_type);
-  bool ValidateTypeDecl(const ast::TypeDecl* named_type) const;
   bool ValidateTextureIntrinsicFunction(const sem::Call* call);
   bool ValidateNoDuplicateDecorations(const ast::DecorationList& decorations);
   // sem::Struct is assumed to have at least one member
@@ -343,6 +341,10 @@
   /// Allocate constant IDs for pipeline-overridable constants.
   void AllocateOverridableConstantIds();
 
+  /// Set the shadowing information on variable declarations.
+  /// @note this method must only be called after all semantic nodes are built.
+  void SetShadows();
+
   /// @returns the resolved type of the ast::Expression `expr`
   /// @param expr the expression
   sem::Type* TypeOf(const ast::Expression* expr);
@@ -410,14 +412,28 @@
   template <typename SEM = sem::Info::InferFromAST,
             typename AST_OR_TYPE = CastableBase>
   auto* Sem(const AST_OR_TYPE* ast) {
-    auto* sem = builder_->Sem().Get<SEM>(ast);
+    using T = sem::Info::GetResultType<SEM, AST_OR_TYPE>;
+    auto* sem = builder_->Sem().Get(ast);
     if (!sem) {
       TINT_ICE(Resolver, diagnostics_)
           << "AST node '" << ast->TypeInfo().name << "' had no semantic info\n"
           << "At: " << ast->source << "\n"
           << "Pointer: " << ast;
     }
-    return sem;
+    return const_cast<T*>(As<T>(sem));
+  }
+
+  /// @returns true if the symbol is the name of an intrinsic (builtin)
+  /// function.
+  bool IsIntrinsic(Symbol) const;
+
+  /// @returns the resolved symbol (function, type or variable) for the given
+  /// ast::Identifier or ast::TypeName cast to the given semantic type.
+  template <typename SEM = sem::Node>
+  SEM* ResolvedSymbol(const ast::Node* node) {
+    auto* resolved = utils::Lookup(dependencies_.resolved_symbols, node);
+    return resolved ? const_cast<SEM*>(builder_->Sem().Get<SEM>(resolved))
+                    : nullptr;
   }
 
   struct TypeConversionSig {
@@ -456,12 +472,8 @@
   diag::List& diagnostics_;
   std::unique_ptr<IntrinsicTable> const intrinsic_table_;
   DependencyGraph dependencies_;
-  ScopeStack<sem::Variable*> variable_stack_;
-  std::unordered_map<Symbol, sem::Function*> symbol_to_function_;
   std::vector<sem::Function*> entry_points_;
   std::unordered_map<const sem::Type*, const Source&> atomic_composite_info_;
-  std::unordered_map<Symbol, TypeDeclInfo> named_type_info_;
-
   std::unordered_set<const ast::Node*> marked_;
   std::unordered_map<uint32_t, const sem::Variable*> constant_ids_;
   std::unordered_map<TypeConversionSig,
diff --git a/src/resolver/resolver_validation.cc b/src/resolver/resolver_validation.cc
index 7b648f1..a5b59f1 100644
--- a/src/resolver/resolver_validation.cc
+++ b/src/resolver/resolver_validation.cc
@@ -942,11 +942,6 @@
 
 bool Resolver::ValidateFunction(const sem::Function* func) {
   auto* decl = func->Declaration();
-  if (!ValidateNoDuplicateDefinition(decl->symbol, decl->source,
-                                     /* check_global_scope_only */ true)) {
-    return false;
-  }
-
   auto workgroup_deco_count = 0;
   for (auto* deco : decl->decorations) {
     if (deco->Is<ast::WorkgroupDecoration>()) {
@@ -1444,7 +1439,7 @@
     if (param_type->Is<sem::Pointer>()) {
       auto is_valid = false;
       if (auto* ident_expr = arg_expr->As<ast::IdentifierExpression>()) {
-        auto* var = variable_stack_.Get(ident_expr->symbol);
+        auto* var = ResolvedSymbol<sem::Variable>(ident_expr);
         if (!var) {
           TINT_ICE(Resolver, diagnostics_) << "failed to resolve identifier";
           return false;
@@ -1456,7 +1451,7 @@
         if (unary->op == ast::UnaryOp::kAddressOf) {
           if (auto* ident_unary =
                   unary->expr->As<ast::IdentifierExpression>()) {
-            auto* var = variable_stack_.Get(ident_unary->symbol);
+            auto* var = ResolvedSymbol<sem::Variable>(ident_unary);
             if (!var) {
               TINT_ICE(Resolver, diagnostics_)
                   << "failed to resolve identifier";
@@ -1755,23 +1750,6 @@
   return true;
 }
 
-bool Resolver::ValidateTypeDecl(const ast::TypeDecl* named_type) const {
-  auto iter = named_type_info_.find(named_type->name);
-  if (iter == named_type_info_.end()) {
-    TINT_ICE(Resolver, diagnostics_)
-        << "ValidateTypeDecl called() before TypeDecl()";
-  }
-  if (iter->second.ast != named_type) {
-    AddError("type with the name '" +
-                 builder_->Symbols().NameFor(named_type->name) +
-                 "' was already declared",
-             named_type->source);
-    AddNote("first declared here", iter->second.ast->source);
-    return false;
-  }
-  return true;
-}
-
 bool Resolver::ValidatePipelineStages() {
   auto check_workgroup_storage = [&](const sem::Function* func,
                                      const sem::Function* entry_point) {
@@ -2338,22 +2316,21 @@
   // https://gpuweb.github.io/gpuweb/wgsl/#assignment-statement
   auto const* lhs_ty = TypeOf(a->lhs);
 
-  if (auto* ident = a->lhs->As<ast::IdentifierExpression>()) {
-    if (auto* var = variable_stack_.Get(ident->symbol)) {
-      if (var->Is<sem::Parameter>()) {
-        AddError("cannot assign to function parameter", a->lhs->source);
-        AddNote("'" + builder_->Symbols().NameFor(ident->symbol) +
-                    "' is declared here:",
-                var->Declaration()->source);
-        return false;
-      }
-      if (var->Declaration()->is_const) {
-        AddError("cannot assign to const", a->lhs->source);
-        AddNote("'" + builder_->Symbols().NameFor(ident->symbol) +
-                    "' is declared here:",
-                var->Declaration()->source);
-        return false;
-      }
+  if (auto* var = ResolvedSymbol<sem::Variable>(a->lhs)) {
+    auto* decl = var->Declaration();
+    if (var->Is<sem::Parameter>()) {
+      AddError("cannot assign to function parameter", a->lhs->source);
+      AddNote("'" + builder_->Symbols().NameFor(decl->symbol) +
+                  "' is declared here:",
+              decl->source);
+      return false;
+    }
+    if (decl->is_const) {
+      AddError("cannot assign to const", a->lhs->source);
+      AddNote("'" + builder_->Symbols().NameFor(decl->symbol) +
+                  "' is declared here:",
+              decl->source);
+      return false;
     }
   }
 
@@ -2388,36 +2365,6 @@
   return true;
 }
 
-bool Resolver::ValidateNoDuplicateDefinition(Symbol sym,
-                                             const Source& source,
-                                             bool check_global_scope_only) {
-  if (check_global_scope_only) {
-    if (auto* var = variable_stack_.Get(sym)) {
-      if (var->Is<sem::GlobalVariable>()) {
-        AddError("redefinition of '" + builder_->Symbols().NameFor(sym) + "'",
-                 source);
-        AddNote("previous definition is here", var->Declaration()->source);
-        return false;
-      }
-    }
-    auto it = symbol_to_function_.find(sym);
-    if (it != symbol_to_function_.end()) {
-      AddError("redefinition of '" + builder_->Symbols().NameFor(sym) + "'",
-               source);
-      AddNote("previous definition is here", it->second->Declaration()->source);
-      return false;
-    }
-  } else {
-    if (auto* var = variable_stack_.Get(sym)) {
-      AddError("redefinition of '" + builder_->Symbols().NameFor(sym) + "'",
-               source);
-      AddNote("previous definition is here", var->Declaration()->source);
-      return false;
-    }
-  }
-  return true;
-}
-
 bool Resolver::ValidateNoDuplicateDecorations(
     const ast::DecorationList& decorations) {
   std::unordered_map<const TypeInfo*, Source> seen;
diff --git a/src/resolver/validation_test.cc b/src/resolver/validation_test.cc
index 790fde5..96557d9 100644
--- a/src/resolver/validation_test.cc
+++ b/src/resolver/validation_test.cc
@@ -150,25 +150,28 @@
 
 TEST_F(ResolverValidationTest, Expr_DontCall_Function) {
   Func("func", {}, ty.void_(), {}, {});
-  auto* ident = create<ast::IdentifierExpression>(
-      Source{{Source::Location{3, 3}, Source::Location{3, 8}}},
-      Symbols().Register("func"));
-  WrapInFunction(ident);
+  WrapInFunction(Expr(Source{{{3, 3}, {3, 8}}}, "func"));
 
   EXPECT_FALSE(r()->Resolve());
   EXPECT_EQ(r()->error(), "3:8 error: missing '(' for function call");
 }
 
 TEST_F(ResolverValidationTest, Expr_DontCall_Intrinsic) {
-  auto* ident = create<ast::IdentifierExpression>(
-      Source{{Source::Location{3, 3}, Source::Location{3, 8}}},
-      Symbols().Register("round"));
-  WrapInFunction(ident);
+  WrapInFunction(Expr(Source{{{3, 3}, {3, 8}}}, "round"));
 
   EXPECT_FALSE(r()->Resolve());
   EXPECT_EQ(r()->error(), "3:8 error: missing '(' for intrinsic call");
 }
 
+TEST_F(ResolverValidationTest, Expr_DontCall_Type) {
+  Alias("T", ty.u32());
+  WrapInFunction(Expr(Source{{{3, 3}, {3, 8}}}, "T"));
+
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(),
+            "3:8 error: missing '(' for type constructor or cast");
+}
+
 TEST_F(ResolverValidationTest,
        AssignmentStmt_InvalidLHS_IntrinsicFunctionName) {
   // normalize = 2;
@@ -220,7 +223,7 @@
 
   Func("my_func", ast::VariableList{}, ty.void_(),
        {
-           Assign(Expr(Source{Source::Location{12, 34}}, "global_var"), 3.14f),
+           Assign(Expr(Source{{12, 34}}, "global_var"), 3.14f),
            Return(),
        });
 
@@ -237,7 +240,7 @@
   auto* cond = Expr(true);
   auto* body = Block(Decl(var));
 
-  SetSource(Source{Source::Location{12, 34}});
+  SetSource(Source{{12, 34}});
   auto* lhs = Expr(Source{{12, 34}}, "a");
   auto* rhs = Expr(3.14f);
 
@@ -350,9 +353,7 @@
 TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadChar) {
   Global("my_vec", ty.vec3<f32>(), ast::StorageClass::kPrivate);
 
-  auto* ident = create<ast::IdentifierExpression>(
-      Source{{Source::Location{3, 3}, Source::Location{3, 7}}},
-      Symbols().Register("xyqz"));
+  auto* ident = Expr(Source{{{3, 3}, {3, 7}}}, "xyqz");
 
   auto* mem = MemberAccessor("my_vec", ident);
   WrapInFunction(mem);
@@ -364,9 +365,7 @@
 TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_MixedChars) {
   Global("my_vec", ty.vec4<f32>(), ast::StorageClass::kPrivate);
 
-  auto* ident = create<ast::IdentifierExpression>(
-      Source{{Source::Location{3, 3}, Source::Location{3, 7}}},
-      Symbols().Register("rgyw"));
+  auto* ident = Expr(Source{{{3, 3}, {3, 7}}}, "rgyw");
 
   auto* mem = MemberAccessor("my_vec", ident);
   WrapInFunction(mem);
@@ -380,9 +379,7 @@
 TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadLength) {
   Global("my_vec", ty.vec3<f32>(), ast::StorageClass::kPrivate);
 
-  auto* ident = create<ast::IdentifierExpression>(
-      Source{{Source::Location{3, 3}, Source::Location{3, 8}}},
-      Symbols().Register("zzzzz"));
+  auto* ident = Expr(Source{{{3, 3}, {3, 8}}}, "zzzzz");
   auto* mem = MemberAccessor("my_vec", ident);
   WrapInFunction(mem);
 
@@ -393,8 +390,7 @@
 TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadIndex) {
   Global("my_vec", ty.vec2<f32>(), ast::StorageClass::kPrivate);
 
-  auto* ident = create<ast::IdentifierExpression>(Source{{3, 3}},
-                                                  Symbols().Register("z"));
+  auto* ident = Expr(Source{{3, 3}}, "z");
   auto* mem = MemberAccessor("my_vec", ident);
   WrapInFunction(mem);
 
@@ -406,9 +402,7 @@
   // var param: vec4<f32>
   // let ret: f32 = *(&param).x;
   auto* param = Var("param", ty.vec4<f32>());
-  auto* x = create<ast::IdentifierExpression>(
-      Source{{Source::Location{3, 3}, Source::Location{3, 8}}},
-      Symbols().Register("x"));
+  auto* x = Expr(Source{{{3, 3}, {3, 8}}}, "x");
 
   auto* addressOf_expr = AddressOf(Source{{12, 34}}, param);
   auto* accessor_expr = MemberAccessor(addressOf_expr, x);
@@ -430,9 +424,7 @@
   auto* p =
       Param("p", ty.pointer(ty.vec4<f32>(), ast::StorageClass::kFunction));
   auto* star_p = Deref(p);
-  auto* z = create<ast::IdentifierExpression>(
-      Source{{Source::Location{3, 3}, Source::Location{3, 8}}},
-      Symbols().Register("z"));
+  auto* z = Expr(Source{{{3, 3}, {3, 8}}}, "z");
   auto* accessor_expr = MemberAccessor(star_p, z);
   auto* x = Var("x", ty.f32(), accessor_expr);
   Func("func", {p}, ty.f32(), {Decl(x), Return(x)});
@@ -446,9 +438,7 @@
   // }
   auto* p =
       Param("p", ty.pointer(ty.vec4<f32>(), ast::StorageClass::kFunction));
-  auto* z = create<ast::IdentifierExpression>(
-      Source{{Source::Location{3, 3}, Source::Location{3, 8}}},
-      Symbols().Register("z"));
+  auto* z = Expr(Source{{{3, 3}, {3, 8}}}, "z");
   auto* accessor_expr = MemberAccessor(p, z);
   auto* star_p = Deref(accessor_expr);
   auto* x = Var("x", ty.f32(), star_p);
@@ -473,7 +463,7 @@
   //     }
   // }
 
-  auto error_loc = Source{Source::Location{12, 34}};
+  auto error_loc = Source{{12, 34}};
   auto* body =
       Block(create<ast::ContinueStatement>(),
             Decl(error_loc, Var("z", ty.i32(), ast::StorageClass::kNone)),
@@ -524,9 +514,9 @@
   //     }
   // }
 
-  auto cont_loc = Source{Source::Location{12, 34}};
-  auto decl_loc = Source{Source::Location{56, 78}};
-  auto ref_loc = Source{Source::Location{90, 12}};
+  auto cont_loc = Source{{12, 34}};
+  auto decl_loc = Source{{56, 78}};
+  auto ref_loc = Source{{90, 12}};
   auto* body =
       Block(If(Expr(true), Block(create<ast::ContinueStatement>(cont_loc))),
             Decl(Var(decl_loc, "z", ty.i32(), ast::StorageClass::kNone)));
@@ -556,9 +546,9 @@
   //     }
   // }
 
-  auto cont_loc = Source{Source::Location{12, 34}};
-  auto decl_loc = Source{Source::Location{56, 78}};
-  auto ref_loc = Source{Source::Location{90, 12}};
+  auto cont_loc = Source{{12, 34}};
+  auto decl_loc = Source{{56, 78}};
+  auto ref_loc = Source{{90, 12}};
   auto* body =
       Block(If(Expr(true), Block(create<ast::ContinueStatement>(cont_loc))),
             Decl(Var(decl_loc, "z", ty.i32(), ast::StorageClass::kNone)));
@@ -590,9 +580,9 @@
   //     }
   // }
 
-  auto cont_loc = Source{Source::Location{12, 34}};
-  auto decl_loc = Source{Source::Location{56, 78}};
-  auto ref_loc = Source{Source::Location{90, 12}};
+  auto cont_loc = Source{{12, 34}};
+  auto decl_loc = Source{{56, 78}};
+  auto ref_loc = Source{{90, 12}};
   auto* body =
       Block(If(Expr(true), Block(create<ast::ContinueStatement>(cont_loc))),
             Decl(Var(decl_loc, "z", ty.i32(), ast::StorageClass::kNone)));
@@ -623,9 +613,9 @@
   //     }
   // }
 
-  auto cont_loc = Source{Source::Location{12, 34}};
-  auto decl_loc = Source{Source::Location{56, 78}};
-  auto ref_loc = Source{Source::Location{90, 12}};
+  auto cont_loc = Source{{12, 34}};
+  auto decl_loc = Source{{56, 78}};
+  auto ref_loc = Source{{90, 12}};
   auto* body =
       Block(If(Expr(true), Block(create<ast::ContinueStatement>(cont_loc))),
             Decl(Var(decl_loc, "z", ty.i32(), ast::StorageClass::kNone)));
@@ -724,7 +714,7 @@
   //     }
   // }
 
-  auto error_loc = Source{Source::Location{12, 34}};
+  auto error_loc = Source{{12, 34}};
   auto* body = Block(Decl(Var("z", ty.i32(), ast::StorageClass::kNone)),
                      create<ast::ContinueStatement>());
   auto* continuing = Block(Assign(Expr(error_loc, "z"), 2));
diff --git a/src/resolver/var_let_test.cc b/src/resolver/var_let_test.cc
index 1b45fac..57d7a14 100644
--- a/src/resolver/var_let_test.cc
+++ b/src/resolver/var_let_test.cc
@@ -58,7 +58,7 @@
            Decl(a),
        });
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 
   // `var` declarations are always of reference type
   ASSERT_TRUE(TypeOf(i)->Is<sem::Reference>());
@@ -114,7 +114,7 @@
            Decl(p),
        });
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 
   // `let` declarations are always of the storage type
   EXPECT_TRUE(TypeOf(i)->Is<sem::I32>());
@@ -153,7 +153,7 @@
 
   WrapInFunction(function);
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 
   ASSERT_TRUE(TypeOf(function)->Is<sem::Reference>());
   ASSERT_TRUE(TypeOf(private_)->Is<sem::Reference>());
@@ -187,7 +187,7 @@
                              create<ast::GroupDecoration>(0),
                          });
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 
   ASSERT_TRUE(TypeOf(storage)->Is<sem::Reference>());
 
@@ -222,7 +222,7 @@
 
   WrapInFunction(ptr);
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 
   ASSERT_TRUE(TypeOf(expr)->Is<sem::Reference>());
   ASSERT_TRUE(TypeOf(ptr)->Is<sem::Pointer>());
@@ -232,6 +232,384 @@
   EXPECT_EQ(TypeOf(ptr)->As<sem::Pointer>()->Access(), ast::Access::kReadWrite);
 }
 
+TEST_F(ResolverVarLetTest, LocalShadowsAlias) {
+  // type a = i32;
+  //
+  // fn X() {
+  //   var a = false;
+  // }
+  //
+  // fn Y() {
+  //   let a = true;
+  // }
+
+  auto* t = Alias("a", ty.i32());
+  auto* v = Var("a", nullptr, Expr(false));
+  auto* l = Const("a", nullptr, Expr(false));
+  Func("X", {}, ty.void_(), {Decl(v)});
+  Func("Y", {}, ty.void_(), {Decl(l)});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* type_t = Sem().Get(t);
+  auto* local_v = Sem().Get<sem::LocalVariable>(v);
+  auto* local_l = Sem().Get<sem::LocalVariable>(l);
+
+  ASSERT_NE(local_v, nullptr);
+  ASSERT_NE(local_l, nullptr);
+
+  EXPECT_EQ(local_v->Shadows(), type_t);
+  EXPECT_EQ(local_l->Shadows(), type_t);
+}
+
+TEST_F(ResolverVarLetTest, LocalShadowsStruct) {
+  // struct a {
+  //   m : i32;
+  // };
+  //
+  // fn X() {
+  //   var a = true;
+  // }
+  //
+  // fn Y() {
+  //   let a = false;
+  // }
+
+  auto* t = Structure("a", {Member("m", ty.i32())});
+  auto* v = Var("a", nullptr, Expr(false));
+  auto* l = Const("a", nullptr, Expr(false));
+  Func("X", {}, ty.void_(), {Decl(v)});
+  Func("Y", {}, ty.void_(), {Decl(l)});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* type_t = Sem().Get(t);
+  auto* local_v = Sem().Get<sem::LocalVariable>(v);
+  auto* local_l = Sem().Get<sem::LocalVariable>(l);
+
+  ASSERT_NE(local_v, nullptr);
+  ASSERT_NE(local_l, nullptr);
+
+  EXPECT_EQ(local_v->Shadows(), type_t);
+  EXPECT_EQ(local_l->Shadows(), type_t);
+}
+
+TEST_F(ResolverVarLetTest, LocalShadowsFunction) {
+  // fn a() {
+  //   var a = true;
+  // }
+  //
+  // fn b() {
+  //   let b = false;
+  // }
+
+  auto* v = Var("a", nullptr, Expr(false));
+  auto* l = Const("b", nullptr, Expr(false));
+  auto* fa = Func("a", {}, ty.void_(), {Decl(v)});
+  auto* fb = Func("b", {}, ty.void_(), {Decl(l)});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* local_v = Sem().Get<sem::LocalVariable>(v);
+  auto* local_l = Sem().Get<sem::LocalVariable>(l);
+  auto* func_a = Sem().Get(fa);
+  auto* func_b = Sem().Get(fb);
+
+  ASSERT_NE(local_v, nullptr);
+  ASSERT_NE(local_l, nullptr);
+  ASSERT_NE(func_a, nullptr);
+  ASSERT_NE(func_b, nullptr);
+
+  EXPECT_EQ(local_v->Shadows(), func_a);
+  EXPECT_EQ(local_l->Shadows(), func_b);
+}
+
+TEST_F(ResolverVarLetTest, LocalShadowsGlobalVar) {
+  // var<private> a : i32;
+  //
+  // fn X() {
+  //   var a = a;
+  // }
+  //
+  // fn Y() {
+  //   let a = a;
+  // }
+
+  auto* g = Global("a", ty.i32(), ast::StorageClass::kPrivate);
+  auto* v = Var("a", nullptr, Expr("a"));
+  auto* l = Const("a", nullptr, Expr("a"));
+  Func("X", {}, ty.void_(), {Decl(v)});
+  Func("Y", {}, ty.void_(), {Decl(l)});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* global = Sem().Get(g);
+  auto* local_v = Sem().Get<sem::LocalVariable>(v);
+  auto* local_l = Sem().Get<sem::LocalVariable>(l);
+
+  ASSERT_NE(local_v, nullptr);
+  ASSERT_NE(local_l, nullptr);
+
+  EXPECT_EQ(local_v->Shadows(), global);
+  EXPECT_EQ(local_l->Shadows(), global);
+
+  auto* user_v =
+      Sem().Get<sem::VariableUser>(local_v->Declaration()->constructor);
+  auto* user_l =
+      Sem().Get<sem::VariableUser>(local_l->Declaration()->constructor);
+
+  ASSERT_NE(user_v, nullptr);
+  ASSERT_NE(user_l, nullptr);
+
+  EXPECT_EQ(user_v->Variable(), global);
+  EXPECT_EQ(user_l->Variable(), global);
+}
+
+TEST_F(ResolverVarLetTest, LocalShadowsGlobalLet) {
+  // let a : i32 = 1;
+  //
+  // fn X() {
+  //   var a = (a == 123);
+  // }
+  //
+  // fn Y() {
+  //   let a = (a == 321);
+  // }
+
+  auto* g = GlobalConst("a", ty.i32(), Expr(1));
+  auto* v = Var("a", nullptr, Expr("a"));
+  auto* l = Const("a", nullptr, Expr("a"));
+  Func("X", {}, ty.void_(), {Decl(v)});
+  Func("Y", {}, ty.void_(), {Decl(l)});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* global = Sem().Get(g);
+  auto* local_v = Sem().Get<sem::LocalVariable>(v);
+  auto* local_l = Sem().Get<sem::LocalVariable>(l);
+
+  ASSERT_NE(local_v, nullptr);
+  ASSERT_NE(local_l, nullptr);
+
+  EXPECT_EQ(local_v->Shadows(), global);
+  EXPECT_EQ(local_l->Shadows(), global);
+
+  auto* user_v =
+      Sem().Get<sem::VariableUser>(local_v->Declaration()->constructor);
+  auto* user_l =
+      Sem().Get<sem::VariableUser>(local_l->Declaration()->constructor);
+
+  ASSERT_NE(user_v, nullptr);
+  ASSERT_NE(user_l, nullptr);
+
+  EXPECT_EQ(user_v->Variable(), global);
+  EXPECT_EQ(user_l->Variable(), global);
+}
+
+TEST_F(ResolverVarLetTest, LocalShadowsLocalVar) {
+  // fn X() {
+  //   var a : i32;
+  //   {
+  //     var a = a;
+  //   }
+  //   {
+  //     let a = a;
+  //   }
+  // }
+
+  auto* s = Var("a", ty.i32(), Expr(1));
+  auto* v = Var("a", nullptr, Expr("a"));
+  auto* l = Const("a", nullptr, Expr("a"));
+  Func("X", {}, ty.void_(), {Decl(s), Block(Decl(v)), Block(Decl(l))});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* local_s = Sem().Get<sem::LocalVariable>(s);
+  auto* local_v = Sem().Get<sem::LocalVariable>(v);
+  auto* local_l = Sem().Get<sem::LocalVariable>(l);
+
+  ASSERT_NE(local_s, nullptr);
+  ASSERT_NE(local_v, nullptr);
+  ASSERT_NE(local_l, nullptr);
+
+  EXPECT_EQ(local_v->Shadows(), local_s);
+  EXPECT_EQ(local_l->Shadows(), local_s);
+
+  auto* user_v =
+      Sem().Get<sem::VariableUser>(local_v->Declaration()->constructor);
+  auto* user_l =
+      Sem().Get<sem::VariableUser>(local_l->Declaration()->constructor);
+
+  ASSERT_NE(user_v, nullptr);
+  ASSERT_NE(user_l, nullptr);
+
+  EXPECT_EQ(user_v->Variable(), local_s);
+  EXPECT_EQ(user_l->Variable(), local_s);
+}
+
+TEST_F(ResolverVarLetTest, LocalShadowsLocalLet) {
+  // fn X() {
+  //   let a = 1;
+  //   {
+  //     var a = (a == 123);
+  //   }
+  //   {
+  //     let a = (a == 321);
+  //   }
+  // }
+
+  auto* s = Const("a", ty.i32(), Expr(1));
+  auto* v = Var("a", nullptr, Expr("a"));
+  auto* l = Const("a", nullptr, Expr("a"));
+  Func("X", {}, ty.void_(), {Decl(s), Block(Decl(v)), Block(Decl(l))});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* local_s = Sem().Get<sem::LocalVariable>(s);
+  auto* local_v = Sem().Get<sem::LocalVariable>(v);
+  auto* local_l = Sem().Get<sem::LocalVariable>(l);
+
+  ASSERT_NE(local_s, nullptr);
+  ASSERT_NE(local_v, nullptr);
+  ASSERT_NE(local_l, nullptr);
+
+  EXPECT_EQ(local_v->Shadows(), local_s);
+  EXPECT_EQ(local_l->Shadows(), local_s);
+
+  auto* user_v =
+      Sem().Get<sem::VariableUser>(local_v->Declaration()->constructor);
+  auto* user_l =
+      Sem().Get<sem::VariableUser>(local_l->Declaration()->constructor);
+
+  ASSERT_NE(user_v, nullptr);
+  ASSERT_NE(user_l, nullptr);
+
+  EXPECT_EQ(user_v->Variable(), local_s);
+  EXPECT_EQ(user_l->Variable(), local_s);
+}
+
+TEST_F(ResolverVarLetTest, LocalShadowsParam) {
+  // fn F(a : i32) {
+  //   {
+  //     var a = a;
+  //   }
+  //   {
+  //     let a = a;
+  //   }
+  // }
+
+  auto* p = Param("a", ty.i32());
+  auto* v = Var("a", nullptr, Expr("a"));
+  auto* l = Const("a", nullptr, Expr("a"));
+  Func("X", {p}, ty.void_(), {Block(Decl(v)), Block(Decl(l))});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* param = Sem().Get<sem::Parameter>(p);
+  auto* local_v = Sem().Get<sem::LocalVariable>(v);
+  auto* local_l = Sem().Get<sem::LocalVariable>(l);
+
+  ASSERT_NE(param, nullptr);
+  ASSERT_NE(local_v, nullptr);
+  ASSERT_NE(local_l, nullptr);
+
+  EXPECT_EQ(local_v->Shadows(), param);
+  EXPECT_EQ(local_l->Shadows(), param);
+
+  auto* user_v =
+      Sem().Get<sem::VariableUser>(local_v->Declaration()->constructor);
+  auto* user_l =
+      Sem().Get<sem::VariableUser>(local_l->Declaration()->constructor);
+
+  ASSERT_NE(user_v, nullptr);
+  ASSERT_NE(user_l, nullptr);
+
+  EXPECT_EQ(user_v->Variable(), param);
+  EXPECT_EQ(user_l->Variable(), param);
+}
+
+TEST_F(ResolverVarLetTest, ParamShadowsFunction) {
+  // fn a(a : bool) {
+  // }
+
+  auto* p = Param("a", ty.bool_());
+  auto* f = Func("a", {p}, ty.void_(), {});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* func = Sem().Get(f);
+  auto* param = Sem().Get<sem::Parameter>(p);
+
+  ASSERT_NE(func, nullptr);
+  ASSERT_NE(param, nullptr);
+
+  EXPECT_EQ(param->Shadows(), func);
+}
+
+TEST_F(ResolverVarLetTest, ParamShadowsGlobalVar) {
+  // var<private> a : i32;
+  //
+  // fn F(a : bool) {
+  // }
+
+  auto* g = Global("a", ty.i32(), ast::StorageClass::kPrivate);
+  auto* p = Param("a", ty.bool_());
+  Func("F", {p}, ty.void_(), {});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* global = Sem().Get(g);
+  auto* param = Sem().Get<sem::Parameter>(p);
+
+  ASSERT_NE(global, nullptr);
+  ASSERT_NE(param, nullptr);
+
+  EXPECT_EQ(param->Shadows(), global);
+}
+
+TEST_F(ResolverVarLetTest, ParamShadowsGlobalLet) {
+  // let a : i32 = 1;
+  //
+  // fn F(a : bool) {
+  // }
+
+  auto* g = GlobalConst("a", ty.i32(), Expr(1));
+  auto* p = Param("a", ty.bool_());
+  Func("F", {p}, ty.void_(), {});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* global = Sem().Get(g);
+  auto* param = Sem().Get<sem::Parameter>(p);
+
+  ASSERT_NE(global, nullptr);
+  ASSERT_NE(param, nullptr);
+
+  EXPECT_EQ(param->Shadows(), global);
+}
+
+TEST_F(ResolverVarLetTest, ParamShadowsAlias) {
+  // type a = i32;
+  //
+  // fn F(a : a) {
+  // }
+
+  auto* a = Alias("a", ty.i32());
+  auto* p = Param("a", ty.type_name("a"));
+  Func("F", {p}, ty.void_(), {});
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+  auto* alias = Sem().Get(a);
+  auto* param = Sem().Get<sem::Parameter>(p);
+
+  ASSERT_NE(alias, nullptr);
+  ASSERT_NE(param, nullptr);
+
+  EXPECT_EQ(param->Shadows(), alias);
+  EXPECT_EQ(param->Type(), alias);
+}
+
 }  // namespace
 }  // namespace resolver
 }  // namespace tint
diff --git a/src/resolver/var_let_validation_test.cc b/src/resolver/var_let_validation_test.cc
index b3a903e..6326f3c 100644
--- a/src/resolver/var_let_validation_test.cc
+++ b/src/resolver/var_let_validation_test.cc
@@ -175,10 +175,7 @@
   WrapInFunction(Var(Source{{12, 34}}, "v", ty.f32(), ast::StorageClass::kNone,
                      Expr(2.0f)));
 
-  EXPECT_FALSE(r()->Resolve()) << r()->error();
-  EXPECT_EQ(
-      r()->error(),
-      "12:34 error: redefinition of 'v'\nnote: previous definition is here");
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverVarLetValidationTest, VarRedeclaredInInnerBlock) {
@@ -194,10 +191,7 @@
 
   WrapInFunction(outer_body);
 
-  EXPECT_FALSE(r()->Resolve());
-  EXPECT_EQ(
-      r()->error(),
-      "12:34 error: redefinition of 'v'\nnote: previous definition is here");
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverVarLetValidationTest, VarRedeclaredInIfBlock) {
@@ -219,10 +213,7 @@
 
   WrapInFunction(outer_body);
 
-  EXPECT_FALSE(r()->Resolve());
-  EXPECT_EQ(
-      r()->error(),
-      "12:34 error: redefinition of 'v'\nnote: previous definition is here");
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverVarLetValidationTest, InferredPtrStorageAccessMismatch) {
diff --git a/src/sem/type_mappings.h b/src/sem/type_mappings.h
index 97a052b..073ff13 100644
--- a/src/sem/type_mappings.h
+++ b/src/sem/type_mappings.h
@@ -25,6 +25,7 @@
 class Expression;
 class Function;
 class MemberAccessorExpression;
+class Node;
 class Statement;
 class Struct;
 class StructMember;
@@ -56,11 +57,12 @@
   Expression* operator()(ast::Expression*);
   Function* operator()(ast::Function*);
   MemberAccessorExpression* operator()(ast::MemberAccessorExpression*);
+  Node* operator()(ast::Node*);
   Statement* operator()(ast::Statement*);
+  Struct* operator()(ast::Struct*);
   StructMember* operator()(ast::StructMember*);
   Type* operator()(ast::Type*);
   Type* operator()(ast::TypeDecl*);
-  Struct* operator()(ast::Struct*);
   Variable* operator()(ast::Variable*);
   //! @endcond
 };
diff --git a/src/sem/variable.cc b/src/sem/variable.cc
index 8b5fdd9..364a3cb 100644
--- a/src/sem/variable.cc
+++ b/src/sem/variable.cc
@@ -45,12 +45,10 @@
                              const sem::Type* type,
                              ast::StorageClass storage_class,
                              ast::Access access,
+                             const sem::Statement* statement,
                              Constant constant_value)
-    : Base(declaration,
-           type,
-           storage_class,
-           access,
-           std::move(constant_value)) {}
+    : Base(declaration, type, storage_class, access, std::move(constant_value)),
+      statement_(statement) {}
 
 LocalVariable::~LocalVariable() = default;
 
diff --git a/src/sem/variable.h b/src/sem/variable.h
index 78e208a..a389eed 100644
--- a/src/sem/variable.h
+++ b/src/sem/variable.h
@@ -95,15 +95,31 @@
   /// @param type the variable type
   /// @param storage_class the variable storage class
   /// @param access the variable access control type
+  /// @param statement the statement that declared this local variable
   /// @param constant_value the constant value for the variable. May be invalid
   LocalVariable(const ast::Variable* declaration,
                 const sem::Type* type,
                 ast::StorageClass storage_class,
                 ast::Access access,
+                const sem::Statement* statement,
                 Constant constant_value);
 
   /// Destructor
   ~LocalVariable() override;
+
+  /// @returns the statement that declares this local variable
+  const sem::Statement* Statement() const { return statement_; }
+
+  /// @returns the Type, Function or Variable that this local variable shadows
+  const sem::Node* Shadows() const { return shadows_; }
+
+  /// Sets the Type, Function or Variable that this local variable shadows
+  /// @param shadows the Type, Function or Variable that this variable shadows
+  void SetShadows(const sem::Node* shadows) { shadows_ = shadows; }
+
+ private:
+  const sem::Statement* const statement_;
+  const sem::Node* shadows_ = nullptr;
 };
 
 /// GlobalVariable is a module-scope variable
@@ -185,10 +201,18 @@
   /// @param owner the CallTarget owner of this parameter
   void SetOwner(CallTarget const* owner) { owner_ = owner; }
 
+  /// @returns the Type, Function or Variable that this local variable shadows
+  const sem::Node* Shadows() const { return shadows_; }
+
+  /// Sets the Type, Function or Variable that this local variable shadows
+  /// @param shadows the Type, Function or Variable that this variable shadows
+  void SetShadows(const sem::Node* shadows) { shadows_ = shadows; }
+
  private:
   const uint32_t index_;
   const ParameterUsage usage_;
   CallTarget const* owner_ = nullptr;
+  const sem::Node* shadows_ = nullptr;
 };
 
 /// ParameterList is a list of Parameter
diff --git a/src/transform/array_length_from_uniform_test.cc b/src/transform/array_length_from_uniform_test.cc
index 13aca0ae..c316412 100644
--- a/src/transform/array_length_from_uniform_test.cc
+++ b/src/transform/array_length_from_uniform_test.cc
@@ -18,6 +18,7 @@
 
 #include "src/transform/simplify_pointers.h"
 #include "src/transform/test_helper.h"
+#include "src/transform/unshadow.h"
 
 namespace tint {
 namespace transform {
@@ -32,7 +33,7 @@
       "error: missing transform data for "
       "tint::transform::ArrayLengthFromUniform";
 
-  auto got = Run<SimplifyPointers, ArrayLengthFromUniform>(src);
+  auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -93,7 +94,7 @@
   DataMap data;
   data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
 
-  auto got = Run<SimplifyPointers, ArrayLengthFromUniform>(src, data);
+  auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
 
   EXPECT_EQ(expect, str(got));
   EXPECT_EQ(std::unordered_set<uint32_t>({0}),
@@ -146,7 +147,7 @@
   DataMap data;
   data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
 
-  auto got = Run<SimplifyPointers, ArrayLengthFromUniform>(src, data);
+  auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
 
   EXPECT_EQ(expect, str(got));
   EXPECT_EQ(std::unordered_set<uint32_t>({0}),
@@ -267,7 +268,7 @@
   DataMap data;
   data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
 
-  auto got = Run<SimplifyPointers, ArrayLengthFromUniform>(src, data);
+  auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
 
   EXPECT_EQ(expect, str(got));
   EXPECT_EQ(std::unordered_set<uint32_t>({0, 1, 2, 3, 4}),
@@ -382,7 +383,7 @@
   DataMap data;
   data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
 
-  auto got = Run<SimplifyPointers, ArrayLengthFromUniform>(src, data);
+  auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
 
   EXPECT_EQ(expect, str(got));
   EXPECT_EQ(std::unordered_set<uint32_t>({0, 2}),
@@ -411,7 +412,7 @@
   DataMap data;
   data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
 
-  auto got = Run<SimplifyPointers, ArrayLengthFromUniform>(src, data);
+  auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
 
   EXPECT_EQ(src, str(got));
   EXPECT_EQ(std::unordered_set<uint32_t>(),
@@ -482,7 +483,7 @@
   DataMap data;
   data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
 
-  auto got = Run<SimplifyPointers, ArrayLengthFromUniform>(src, data);
+  auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
 
   EXPECT_EQ(expect, str(got));
   EXPECT_EQ(std::unordered_set<uint32_t>({0}),
diff --git a/src/transform/calculate_array_length.cc b/src/transform/calculate_array_length.cc
index 0507685..e06b959 100644
--- a/src/transform/calculate_array_length.cc
+++ b/src/transform/calculate_array_length.cc
@@ -25,8 +25,9 @@
 #include "src/sem/statement.h"
 #include "src/sem/struct.h"
 #include "src/sem/variable.h"
-#include "src/utils/map.h"
+#include "src/transform/simplify_pointers.h"
 #include "src/utils/hash.h"
+#include "src/utils/map.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::transform::CalculateArrayLength);
 TINT_INSTANTIATE_TYPEINFO(
@@ -72,6 +73,9 @@
 
 void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) {
   auto& sem = ctx.src->Sem();
+  if (!Requires<SimplifyPointers>(ctx)) {
+    return;
+  }
 
   // get_buffer_size_intrinsic() emits the function decorated with
   // BufferSizeIntrinsic that is transformed by the HLSL writer into a call to
diff --git a/src/transform/calculate_array_length_test.cc b/src/transform/calculate_array_length_test.cc
index 1e72a40..897213c 100644
--- a/src/transform/calculate_array_length_test.cc
+++ b/src/transform/calculate_array_length_test.cc
@@ -14,7 +14,9 @@
 
 #include "src/transform/calculate_array_length.h"
 
+#include "src/transform/simplify_pointers.h"
 #include "src/transform/test_helper.h"
+#include "src/transform/unshadow.h"
 
 namespace tint {
 namespace transform {
@@ -22,6 +24,18 @@
 
 using CalculateArrayLengthTest = TransformTest;
 
+TEST_F(CalculateArrayLengthTest, Error_MissingCalculateArrayLength) {
+  auto* src = "";
+
+  auto* expect =
+      "error: tint::transform::CalculateArrayLength depends on "
+      "tint::transform::SimplifyPointers but the dependency was not run";
+
+  auto got = Run<CalculateArrayLength>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(CalculateArrayLengthTest, Basic) {
   auto* src = R"(
 [[block]]
@@ -59,7 +73,7 @@
 }
 )";
 
-  auto got = Run<CalculateArrayLength>(src);
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -105,7 +119,7 @@
 }
 )";
 
-  auto got = Run<CalculateArrayLength>(src);
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -149,7 +163,7 @@
 }
 )";
 
-  auto got = Run<CalculateArrayLength>(src);
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -206,7 +220,7 @@
 }
 )";
 
-  auto got = Run<CalculateArrayLength>(src);
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -274,7 +288,63 @@
 }
 )";
 
-  auto got = Run<CalculateArrayLength>(src);
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(CalculateArrayLengthTest, Shadowing) {
+  auto* src = R"(
+[[block]]
+struct SB {
+  x : i32;
+  arr : array<i32>;
+};
+
+[[group(0), binding(0)]] var<storage, read> a : SB;
+[[group(0), binding(1)]] var<storage, read> b : SB;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = &a;
+  var a : u32 = arrayLength(&a.arr);
+  {
+    var b : u32 = arrayLength(&((*x).arr));
+  }
+}
+)";
+
+  auto* expect =
+      R"(
+[[block]]
+struct SB {
+  x : i32;
+  arr : array<i32>;
+};
+
+[[internal(intrinsic_buffer_size)]]
+fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, result : ptr<function, u32>)
+
+[[group(0), binding(0)]] var<storage, read> a : SB;
+
+[[group(0), binding(1)]] var<storage, read> b : SB;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var tint_symbol_1 : u32 = 0u;
+  tint_symbol(a, &(tint_symbol_1));
+  let tint_symbol_2 : u32 = ((tint_symbol_1 - 4u) / 4u);
+  var a_1 : u32 = tint_symbol_2;
+  {
+    var tint_symbol_3 : u32 = 0u;
+    tint_symbol(a, &(tint_symbol_3));
+    let tint_symbol_4 : u32 = ((tint_symbol_3 - 4u) / 4u);
+    var b_1 : u32 = tint_symbol_4;
+  }
+}
+)";
+
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
 
   EXPECT_EQ(expect, str(got));
 }
diff --git a/src/transform/canonicalize_entry_point_io.cc b/src/transform/canonicalize_entry_point_io.cc
index af54484..c78ac19 100644
--- a/src/transform/canonicalize_entry_point_io.cc
+++ b/src/transform/canonicalize_entry_point_io.cc
@@ -23,6 +23,7 @@
 #include "src/ast/disable_validation_decoration.h"
 #include "src/program_builder.h"
 #include "src/sem/function.h"
+#include "src/transform/unshadow.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::transform::CanonicalizeEntryPointIO);
 TINT_INSTANTIATE_TYPEINFO(tint::transform::CanonicalizeEntryPointIO::Config);
@@ -550,6 +551,10 @@
 void CanonicalizeEntryPointIO::Run(CloneContext& ctx,
                                    const DataMap& inputs,
                                    DataMap&) {
+  if (!Requires<Unshadow>(ctx)) {
+    return;
+  }
+
   auto* cfg = inputs.Get<Config>();
   if (cfg == nullptr) {
     ctx.dst->Diagnostics().add_error(
diff --git a/src/transform/canonicalize_entry_point_io_test.cc b/src/transform/canonicalize_entry_point_io_test.cc
index a011e85..3ea25c1 100644
--- a/src/transform/canonicalize_entry_point_io_test.cc
+++ b/src/transform/canonicalize_entry_point_io_test.cc
@@ -15,6 +15,7 @@
 #include "src/transform/canonicalize_entry_point_io.h"
 
 #include "src/transform/test_helper.h"
+#include "src/transform/unshadow.h"
 
 namespace tint {
 namespace transform {
@@ -22,6 +23,18 @@
 
 using CanonicalizeEntryPointIOTest = TransformTest;
 
+TEST_F(CanonicalizeEntryPointIOTest, Error_MissingUnshadow) {
+  auto* src = "";
+
+  auto* expect =
+      "error: tint::transform::CanonicalizeEntryPointIO depends on "
+      "tint::transform::Unshadow but the dependency was not run";
+
+  auto got = Run<CanonicalizeEntryPointIO>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(CanonicalizeEntryPointIOTest, Error_MissingTransformData) {
   auto* src = "";
 
@@ -29,7 +42,7 @@
       "error: missing transform data for "
       "tint::transform::CanonicalizeEntryPointIO";
 
-  auto got = Run<CanonicalizeEntryPointIO>(src);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -52,7 +65,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -87,7 +100,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kSpirv);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -123,7 +136,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -161,7 +174,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -197,7 +210,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -251,7 +264,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kSpirv);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -306,7 +319,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -363,7 +376,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -393,7 +406,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kSpirv);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -428,7 +441,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -463,7 +476,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -519,7 +532,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kSpirv);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -580,7 +593,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -641,7 +654,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -709,7 +722,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kSpirv);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -783,7 +796,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -857,7 +870,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -925,7 +938,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1012,7 +1025,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1104,7 +1117,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1235,7 +1248,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kSpirv);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1300,7 +1313,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1370,7 +1383,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1475,7 +1488,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1505,7 +1518,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1538,7 +1551,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1573,7 +1586,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1608,7 +1621,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1646,7 +1659,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1699,7 +1712,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1751,7 +1764,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1841,7 +1854,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03u);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1892,7 +1905,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0x03);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1926,7 +1939,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kSpirv, 0xFFFFFFFF, true);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -1964,7 +1977,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0xFFFFFFFF, true);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -2005,7 +2018,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kSpirv, 0xFFFFFFFF, true);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -2051,7 +2064,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0xFFFFFFFF, true);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -2129,7 +2142,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kSpirv, 0xFFFFFFFF, true);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -2205,7 +2218,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kMsl, 0xFFFFFFFF, true);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -2281,7 +2294,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl, 0xFFFFFFFF, true);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -2317,7 +2330,7 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kSpirv);
-  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO>(src, data);
 
   EXPECT_EQ(expect, str(got));
 }
diff --git a/src/transform/decompose_strided_matrix_test.cc b/src/transform/decompose_strided_matrix_test.cc
index 714e8a1..b7a22e4 100644
--- a/src/transform/decompose_strided_matrix_test.cc
+++ b/src/transform/decompose_strided_matrix_test.cc
@@ -22,6 +22,7 @@
 #include "src/program_builder.h"
 #include "src/transform/simplify_pointers.h"
 #include "src/transform/test_helper.h"
+#include "src/transform/unshadow.h"
 
 namespace tint {
 namespace transform {
@@ -34,7 +35,7 @@
   auto* src = R"()";
   auto* expect = src;
 
-  auto got = Run<SimplifyPointers, DecomposeStridedMatrix>(src);
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -109,8 +110,8 @@
 }
 )";
 
-  auto got =
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -171,8 +172,8 @@
 }
 )";
 
-  auto got =
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -234,8 +235,8 @@
 }
 )";
 
-  auto got =
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -300,8 +301,8 @@
 }
 )";
 
-  auto got =
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -362,8 +363,8 @@
 }
 )";
 
-  auto got =
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -429,8 +430,8 @@
 }
 )";
 
-  auto got =
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -491,8 +492,8 @@
 }
 )";
 
-  auto got =
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -580,8 +581,8 @@
 }
 )";
 
-  auto got =
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -637,8 +638,8 @@
 }
 )";
 
-  auto got =
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -695,8 +696,8 @@
 }
 )";
 
-  auto got =
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
diff --git a/src/transform/fold_trivial_single_use_lets_test.cc b/src/transform/fold_trivial_single_use_lets_test.cc
index b571ce8..42aae29 100644
--- a/src/transform/fold_trivial_single_use_lets_test.cc
+++ b/src/transform/fold_trivial_single_use_lets_test.cc
@@ -35,13 +35,13 @@
   auto* src = R"(
 fn f() {
   let x = 1;
-  ignore(x);
+  _ = x;
 }
 )";
 
   auto* expect = R"(
 fn f() {
-  ignore(1);
+  _ = 1;
 }
 )";
 
@@ -56,13 +56,13 @@
   let x = 1;
   let y = 2;
   let z = 3;
-  ignore(x + y + z);
+  _ = x + y + z;
 }
 )";
 
   auto* expect = R"(
 fn f() {
-  ignore(((1 + 2) + 3));
+  _ = ((1 + 2) + 3);
 }
 )";
 
@@ -77,13 +77,13 @@
   let x = 1;
   let y = x;
   let z = y;
-  ignore(z);
+  _ = z;
 }
 )";
 
   auto* expect = R"(
 fn f() {
-  ignore(1);
+  _ = 1;
 }
 )";
 
@@ -101,7 +101,7 @@
 fn f() {
   let x = 1;
   let y = function_with_posssible_side_effect();
-  ignore((x + y));
+  _ = (x + y);
 }
 )";
 
@@ -117,7 +117,7 @@
 fn f() {
   let x = 1;
   {
-    ignore(x);
+    _ = x;
   }
 }
 )";
@@ -133,7 +133,26 @@
   auto* src = R"(
 fn f() {
   let x = 1;
-  ignore((x + x));
+  _ = (x + x);
+}
+)";
+
+  auto* expect = src;
+
+  auto got = Run<FoldTrivialSingleUseLets>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(FoldTrivialSingleUseLetsTest, NoFold_Shadowing) {
+  auto* src = R"(
+fn f() {
+  var y = 1;
+  let x = y;
+  {
+    let y = false;
+    _ = (x + x);
+  }
 }
 )";
 
diff --git a/src/transform/glsl.cc b/src/transform/glsl.cc
index 4b4512d..1000234 100644
--- a/src/transform/glsl.cc
+++ b/src/transform/glsl.cc
@@ -29,6 +29,7 @@
 #include "src/transform/remove_phonies.h"
 #include "src/transform/simplify_pointers.h"
 #include "src/transform/single_entry_point.h"
+#include "src/transform/unshadow.h"
 #include "src/transform/zero_init_workgroup_memory.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::transform::Glsl);
@@ -46,6 +47,8 @@
 
   auto* cfg = inputs.Get<Config>();
 
+  manager.Add<Unshadow>();
+
   // Attempt to convert `loop`s into for-loops. This is to try and massage the
   // output into something that will not cause FXC to choke or misbehave.
   manager.Add<FoldTrivialSingleUseLets>();
diff --git a/src/transform/num_workgroups_from_uniform_test.cc b/src/transform/num_workgroups_from_uniform_test.cc
index 16973a4..1d71941 100644
--- a/src/transform/num_workgroups_from_uniform_test.cc
+++ b/src/transform/num_workgroups_from_uniform_test.cc
@@ -18,6 +18,7 @@
 
 #include "src/transform/canonicalize_entry_point_io.h"
 #include "src/transform/test_helper.h"
+#include "src/transform/unshadow.h"
 
 namespace tint {
 namespace transform {
@@ -35,7 +36,8 @@
   DataMap data;
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
-  auto got = Run<CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(
+      src, data);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -87,7 +89,8 @@
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
   data.Add<NumWorkgroupsFromUniform::Config>(sem::BindingPoint{0, 30u});
-  auto got = Run<CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(
+      src, data);
   EXPECT_EQ(expect, str(got));
 }
 
@@ -133,7 +136,8 @@
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
   data.Add<NumWorkgroupsFromUniform::Config>(sem::BindingPoint{0, 30u});
-  auto got = Run<CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(
+      src, data);
   EXPECT_EQ(expect, str(got));
 }
 
@@ -190,7 +194,8 @@
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
   data.Add<NumWorkgroupsFromUniform::Config>(sem::BindingPoint{0, 30u});
-  auto got = Run<CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(
+      src, data);
   EXPECT_EQ(expect, str(got));
 }
 
@@ -291,7 +296,8 @@
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
   data.Add<NumWorkgroupsFromUniform::Config>(sem::BindingPoint{0, 30u});
-  auto got = Run<CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(
+      src, data);
   EXPECT_EQ(expect, str(got));
 }
 
@@ -333,7 +339,8 @@
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
   data.Add<NumWorkgroupsFromUniform::Config>(sem::BindingPoint{0, 30u});
-  auto got = Run<CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(src, data);
+  auto got = Run<Unshadow, CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(
+      src, data);
   EXPECT_EQ(expect, str(got));
 }
 
diff --git a/src/transform/simplify_pointers.cc b/src/transform/simplify_pointers.cc
index 6ab7b6a..b789e87 100644
--- a/src/transform/simplify_pointers.cc
+++ b/src/transform/simplify_pointers.cc
@@ -24,6 +24,7 @@
 #include "src/sem/function.h"
 #include "src/sem/statement.h"
 #include "src/sem/variable.h"
+#include "src/transform/unshadow.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::transform::SimplifyPointers);
 
@@ -231,6 +232,9 @@
 SimplifyPointers::~SimplifyPointers() = default;
 
 void SimplifyPointers::Run(CloneContext& ctx, const DataMap&, DataMap&) {
+  if (!Requires<Unshadow>(ctx)) {
+    return;
+  }
   State(ctx).Run();
 }
 
diff --git a/src/transform/simplify_pointers.h b/src/transform/simplify_pointers.h
index 9f00f9c..1f8cf71 100644
--- a/src/transform/simplify_pointers.h
+++ b/src/transform/simplify_pointers.h
@@ -15,16 +15,14 @@
 #ifndef SRC_TRANSFORM_SIMPLIFY_POINTERS_H_
 #define SRC_TRANSFORM_SIMPLIFY_POINTERS_H_
 
-#include <string>
-#include <unordered_map>
-
 #include "src/transform/transform.h"
 
 namespace tint {
 namespace transform {
 
 /// SimplifyPointers is a Transform that moves all usage of function-scope
-/// `let` statements of a pointer type into their places of usage.
+/// `let` statements of a pointer type into their places of usage, while also
+/// simplifying any chains of address-of or indirections operators.
 ///
 /// Parameters of a pointer type are not adjusted.
 ///
diff --git a/src/transform/simplify_pointers_test.cc b/src/transform/simplify_pointers_test.cc
index b2816ab..d6fde98 100644
--- a/src/transform/simplify_pointers_test.cc
+++ b/src/transform/simplify_pointers_test.cc
@@ -14,11 +14,8 @@
 
 #include "src/transform/simplify_pointers.h"
 
-#include <memory>
-#include <utility>
-#include <vector>
-
 #include "src/transform/test_helper.h"
+#include "src/transform/unshadow.h"
 
 namespace tint {
 namespace transform {
@@ -26,11 +23,23 @@
 
 using SimplifyPointersTest = TransformTest;
 
+TEST_F(SimplifyPointersTest, Error_MissingSimplifyPointers) {
+  auto* src = "";
+
+  auto* expect =
+      "error: tint::transform::SimplifyPointers depends on "
+      "tint::transform::Unshadow but the dependency was not run";
+
+  auto got = Run<SimplifyPointers>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(SimplifyPointersTest, EmptyModule) {
   auto* src = "";
   auto* expect = "";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -51,7 +60,7 @@
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -79,7 +88,7 @@
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -103,7 +112,7 @@
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -126,7 +135,7 @@
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -180,7 +189,7 @@
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -201,7 +210,7 @@
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -229,7 +238,7 @@
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -264,7 +273,7 @@
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -330,34 +339,42 @@
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
 
-// TODO(crbug.com/tint/819): Enable when we support inter-scope shadowing.
-TEST_F(SimplifyPointersTest, DISABLED_ModificationAfterInline) {
+TEST_F(SimplifyPointersTest, ShadowPointer) {
   auto* src = R"(
-fn x(p : ptr<function, i32>) -> i32 {
-  return *p;
-}
+var<private> a : array<i32, 2>;
 
-fn f() {
-  var i : i32 = 1;
-  let p : ptr<function, i32> = &i;
-  if (true) {
-    var i : i32 = 2;
-    x(p);
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = &a;
+  var a : i32 = (*x)[0];
+  {
+    var a : i32 = (*x)[1];
   }
 }
 )";
 
-  auto* expect = R"(<TODO>)";
+  auto* expect = R"(
+var<private> a : array<i32, 2>;
 
-  auto got = Run<SimplifyPointers>(src);
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var a_1 : i32 = a[0];
+  {
+    var a_2 : i32 = a[1];
+  }
+}
+)";
+
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint
diff --git a/src/transform/unshadow.cc b/src/transform/unshadow.cc
new file mode 100644
index 0000000..a7bc66d
--- /dev/null
+++ b/src/transform/unshadow.cc
@@ -0,0 +1,99 @@
+// Copyright 2021 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/transform/unshadow.h"
+
+#include <memory>
+#include <unordered_map>
+#include <utility>
+
+#include "src/program_builder.h"
+#include "src/sem/block_statement.h"
+#include "src/sem/function.h"
+#include "src/sem/statement.h"
+#include "src/sem/variable.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::transform::Unshadow);
+
+namespace tint {
+namespace transform {
+
+/// The PIMPL state for the Unshadow transform
+struct Unshadow::State {
+  /// The clone context
+  CloneContext& ctx;
+
+  /// Constructor
+  /// @param context the clone context
+  explicit State(CloneContext& context) : ctx(context) {}
+
+  /// Performs the transformation
+  void Run() {
+    auto& sem = ctx.src->Sem();
+
+    // Maps a variable to its new name.
+    std::unordered_map<const sem::Variable*, Symbol> renamed_to;
+
+    auto rename = [&](const sem::Variable* var) -> const ast::Variable* {
+      auto* decl = var->Declaration();
+      auto name = ctx.src->Symbols().NameFor(decl->symbol);
+      auto symbol = ctx.dst->Symbols().New(name);
+      renamed_to.emplace(var, symbol);
+
+      auto source = ctx.Clone(decl->source);
+      auto* type = ctx.Clone(decl->type);
+      auto* constructor = ctx.Clone(decl->constructor);
+      auto decorations = ctx.Clone(decl->decorations);
+      return ctx.dst->create<ast::Variable>(
+          source, symbol, decl->declared_storage_class, decl->declared_access,
+          type, decl->is_const, constructor, decorations);
+    };
+
+    ctx.ReplaceAll([&](const ast::Variable* var) -> const ast::Variable* {
+      if (auto* local = sem.Get<sem::LocalVariable>(var)) {
+        if (local->Shadows()) {
+          return rename(local);
+        }
+      }
+      if (auto* param = sem.Get<sem::Parameter>(var)) {
+        if (param->Shadows()) {
+          return rename(param);
+        }
+      }
+      return nullptr;
+    });
+    ctx.ReplaceAll([&](const ast::IdentifierExpression* ident)
+                       -> const tint::ast::IdentifierExpression* {
+      if (auto* user = sem.Get<sem::VariableUser>(ident)) {
+        auto it = renamed_to.find(user->Variable());
+        if (it != renamed_to.end()) {
+          return ctx.dst->Expr(it->second);
+        }
+      }
+      return nullptr;
+    });
+    ctx.Clone();
+  }
+};
+
+Unshadow::Unshadow() = default;
+
+Unshadow::~Unshadow() = default;
+
+void Unshadow::Run(CloneContext& ctx, const DataMap&, DataMap&) {
+  State(ctx).Run();
+}
+
+}  // namespace transform
+}  // namespace tint
diff --git a/src/transform/unshadow.h b/src/transform/unshadow.h
new file mode 100644
index 0000000..7c16bcf
--- /dev/null
+++ b/src/transform/unshadow.h
@@ -0,0 +1,48 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TRANSFORM_UNSHADOW_H_
+#define SRC_TRANSFORM_UNSHADOW_H_
+
+#include "src/transform/transform.h"
+
+namespace tint {
+namespace transform {
+
+/// Unshadow is a Transform that renames any variables that shadow another
+/// variable.
+class Unshadow : public Castable<Unshadow, Transform> {
+ public:
+  /// Constructor
+  Unshadow();
+
+  /// Destructor
+  ~Unshadow() override;
+
+ protected:
+  struct State;
+
+  /// Runs the transform using the CloneContext built for transforming a
+  /// program. Run() is responsible for calling Clone() on the CloneContext.
+  /// @param ctx the CloneContext primed with the input program and
+  /// ProgramBuilder
+  /// @param inputs optional extra transform-specific input data
+  /// @param outputs optional extra transform-specific output data
+  void Run(CloneContext& ctx, const DataMap& inputs, DataMap& outputs) override;
+};
+
+}  // namespace transform
+}  // namespace tint
+
+#endif  // SRC_TRANSFORM_UNSHADOW_H_
diff --git a/src/transform/unshadow_test.cc b/src/transform/unshadow_test.cc
new file mode 100644
index 0000000..2a4cbed
--- /dev/null
+++ b/src/transform/unshadow_test.cc
@@ -0,0 +1,397 @@
+// Copyright 2021 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/transform/unshadow.h"
+
+#include "src/transform/test_helper.h"
+
+namespace tint {
+namespace transform {
+namespace {
+
+using UnshadowTest = TransformTest;
+
+TEST_F(UnshadowTest, EmptyModule) {
+  auto* src = "";
+  auto* expect = "";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, Noop) {
+  auto* src = R"(
+var<private> a : i32;
+
+let b : i32 = 1;
+
+fn F(c : i32) {
+  var d : i32;
+  let e : i32 = 1;
+  {
+    var f : i32;
+    let g : i32 = 1;
+  }
+}
+)";
+
+  auto* expect = src;
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, LocalShadowsAlias) {
+  auto* src = R"(
+type a = i32;
+
+fn X() {
+  var a = false;
+}
+
+fn Y() {
+  let a = true;
+}
+)";
+
+  auto* expect = R"(
+type a = i32;
+
+fn X() {
+  var a_1 = false;
+}
+
+fn Y() {
+  let a_2 = true;
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, LocalShadowsStruct) {
+  auto* src = R"(
+struct a {
+  m : i32;
+};
+
+fn X() {
+  var a = true;
+}
+
+fn Y() {
+  let a = false;
+}
+)";
+
+  auto* expect = R"(
+struct a {
+  m : i32;
+};
+
+fn X() {
+  var a_1 = true;
+}
+
+fn Y() {
+  let a_2 = false;
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, LocalShadowsFunction) {
+  auto* src = R"(
+fn a() {
+  var a = true;
+}
+
+fn b() {
+  let b = false;
+}
+)";
+
+  auto* expect = R"(
+fn a() {
+  var a_1 = true;
+}
+
+fn b() {
+  let b_1 = false;
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, LocalShadowsGlobalVar) {
+  auto* src = R"(
+var<private> a : i32;
+
+fn X() {
+  var a = (a == 123);
+}
+
+fn Y() {
+  let a = (a == 321);
+}
+)";
+
+  auto* expect = R"(
+var<private> a : i32;
+
+fn X() {
+  var a_1 = (a == 123);
+}
+
+fn Y() {
+  let a_2 = (a == 321);
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, LocalShadowsGlobalLet) {
+  auto* src = R"(
+let a : i32 = 1;
+
+fn X() {
+  var a = (a == 123);
+}
+
+fn Y() {
+  let a = (a == 321);
+}
+)";
+
+  auto* expect = R"(
+let a : i32 = 1;
+
+fn X() {
+  var a_1 = (a == 123);
+}
+
+fn Y() {
+  let a_2 = (a == 321);
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, LocalShadowsLocalVar) {
+  auto* src = R"(
+fn X() {
+  var a : i32;
+  {
+    var a = (a == 123);
+  }
+  {
+    let a = (a == 321);
+  }
+}
+)";
+
+  auto* expect = R"(
+fn X() {
+  var a : i32;
+  {
+    var a_1 = (a == 123);
+  }
+  {
+    let a_2 = (a == 321);
+  }
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, LocalShadowsLocalLet) {
+  auto* src = R"(
+fn X() {
+  let a = 1;
+  {
+    var a = (a == 123);
+  }
+  {
+    let a = (a == 321);
+  }
+}
+)";
+
+  auto* expect = R"(
+fn X() {
+  let a = 1;
+  {
+    var a_1 = (a == 123);
+  }
+  {
+    let a_2 = (a == 321);
+  }
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, LocalShadowsParam) {
+  auto* src = R"(
+fn F(a : i32) {
+  {
+    var a = (a == 123);
+  }
+  {
+    let a = (a == 321);
+  }
+}
+)";
+
+  auto* expect = R"(
+fn F(a : i32) {
+  {
+    var a_1 = (a == 123);
+  }
+  {
+    let a_2 = (a == 321);
+  }
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, ParamShadowsFunction) {
+  auto* src = R"(
+fn a(a : i32) {
+  {
+    var a = (a == 123);
+  }
+  {
+    let a = (a == 321);
+  }
+}
+)";
+
+  auto* expect = R"(
+fn a(a_1 : i32) {
+  {
+    var a_2 = (a_1 == 123);
+  }
+  {
+    let a_3 = (a_1 == 321);
+  }
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, ParamShadowsGlobalVar) {
+  auto* src = R"(
+var<private> a : i32;
+
+fn F(a : bool) {
+}
+)";
+
+  auto* expect = R"(
+var<private> a : i32;
+
+fn F(a_1 : bool) {
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, ParamShadowsGlobalLet) {
+  auto* src = R"(
+let a : i32 = 1;
+
+fn F(a : bool) {
+}
+)";
+
+  auto* expect = R"(
+let a : i32 = 1;
+
+fn F(a_1 : bool) {
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(UnshadowTest, ParamShadowsAlias) {
+  auto* src = R"(
+type a = i32;
+
+fn F(a : a) {
+  {
+    var a = (a == 123);
+  }
+  {
+    let a = (a == 321);
+  }
+}
+)";
+
+  auto* expect = R"(
+type a = i32;
+
+fn F(a_1 : a) {
+  {
+    var a_2 = (a_1 == 123);
+  }
+  {
+    let a_3 = (a_1 == 321);
+  }
+}
+)";
+
+  auto got = Run<Unshadow>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+}  // namespace
+}  // namespace transform
+}  // namespace tint
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 3bf5394..1a24ed2 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -58,6 +58,7 @@
 #include "src/transform/promote_initializers_to_const_var.h"
 #include "src/transform/remove_phonies.h"
 #include "src/transform/simplify_pointers.h"
+#include "src/transform/unshadow.h"
 #include "src/transform/zero_init_workgroup_memory.h"
 #include "src/utils/defer.h"
 #include "src/utils/map.h"
@@ -135,6 +136,8 @@
   array_length_from_uniform_cfg.bindpoint_to_size_index =
       array_length_from_uniform.bindpoint_to_size_index;
 
+  manager.Add<transform::Unshadow>();
+
   // Attempt to convert `loop`s into for-loops. This is to try and massage the
   // output into something that will not cause FXC to choke or misbehave.
   manager.Add<transform::FoldTrivialSingleUseLets>();
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index 6054609..077beae 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -66,6 +66,7 @@
 #include "src/transform/promote_initializers_to_const_var.h"
 #include "src/transform/remove_phonies.h"
 #include "src/transform/simplify_pointers.h"
+#include "src/transform/unshadow.h"
 #include "src/transform/vectorize_scalar_matrix_constructors.h"
 #include "src/transform/wrap_arrays_in_structs.h"
 #include "src/transform/zero_init_workgroup_memory.h"
@@ -152,6 +153,8 @@
       transform::CanonicalizeEntryPointIO::ShaderStyle::kMsl, fixed_sample_mask,
       emit_vertex_point_size);
 
+  manager.Add<transform::Unshadow>();
+
   if (!disable_workgroup_init) {
     // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
     // ZeroInitWorkgroupMemory may inject new builtin parameters.
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index b911359..1bbb2b6 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -47,6 +47,7 @@
 #include "src/transform/for_loop_to_loop.h"
 #include "src/transform/manager.h"
 #include "src/transform/simplify_pointers.h"
+#include "src/transform/unshadow.h"
 #include "src/transform/vectorize_scalar_matrix_constructors.h"
 #include "src/transform/zero_init_workgroup_memory.h"
 #include "src/utils/defer.h"
@@ -266,6 +267,7 @@
   transform::Manager manager;
   transform::DataMap data;
 
+  manager.Add<transform::Unshadow>();
   if (!disable_workgroup_init) {
     manager.Add<transform::ZeroInitWorkgroupMemory>();
   }
diff --git a/test.wgsl b/test.wgsl
new file mode 100644
index 0000000..827d200
--- /dev/null
+++ b/test.wgsl
@@ -0,0 +1,15 @@
+[[block]]
+struct SB {
+  x : i32;
+  arr : array<i32>;
+};
+
+[[group(0), binding(0)]] var<storage, read> sb : SB;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var a : u32 = arrayLength(&sb.arr);
+  let p = &sb;
+  let sb2 = p;
+  var b : u32 = arrayLength(&((*sb2).arr));
+}
diff --git a/test/BUILD.gn b/test/BUILD.gn
index 41e1069..a634ea4 100644
--- a/test/BUILD.gn
+++ b/test/BUILD.gn
@@ -327,6 +327,7 @@
     "../src/transform/single_entry_point_test.cc",
     "../src/transform/test_helper.h",
     "../src/transform/transform_test.cc",
+    "../src/transform/unshadow_test.cc",
     "../src/transform/vectorize_scalar_matrix_constructors_test.cc",
     "../src/transform/vertex_pulling_test.cc",
     "../src/transform/wrap_arrays_in_structs_test.cc",
diff --git a/test/bug/tint/825.wgsl.expected.hlsl b/test/bug/tint/825.wgsl.expected.hlsl
index 1e1f290..febe29b 100644
--- a/test/bug/tint/825.wgsl.expected.hlsl
+++ b/test/bug/tint/825.wgsl.expected.hlsl
@@ -7,5 +7,5 @@
   int i = 0;
   int j = 0;
   float2x2 m = float2x2(float2(1.0f, 2.0f), float2(3.0f, 4.0f));
-  const float f = m[i][j];
+  const float f_1 = m[i][j];
 }
diff --git a/test/bug/tint/825.wgsl.expected.msl b/test/bug/tint/825.wgsl.expected.msl
index 58b800a..4630243 100644
--- a/test/bug/tint/825.wgsl.expected.msl
+++ b/test/bug/tint/825.wgsl.expected.msl
@@ -5,6 +5,6 @@
   int i = 0;
   int j = 0;
   float2x2 m = float2x2(float2(1.0f, 2.0f), float2(3.0f, 4.0f));
-  float const f = m[i][j];
+  float const f_1 = m[i][j];
 }
 
diff --git a/test/shadowing/alias/let.wgsl b/test/shadowing/alias/let.wgsl
new file mode 100644
index 0000000..558efe6
--- /dev/null
+++ b/test/shadowing/alias/let.wgsl
@@ -0,0 +1,10 @@
+type a = i32;
+
+fn f() {
+  {
+    let a : a = a();
+    let b = a;
+  }
+  let a : a = a();
+  let b = a;
+}
diff --git a/test/shadowing/alias/let.wgsl.expected.hlsl b/test/shadowing/alias/let.wgsl.expected.hlsl
new file mode 100644
index 0000000..ad5e104
--- /dev/null
+++ b/test/shadowing/alias/let.wgsl.expected.hlsl
@@ -0,0 +1,13 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+void f() {
+  {
+    const int a_1 = 0;
+    const int b = a_1;
+  }
+  const int a_2 = 0;
+  const int b = a_2;
+}
diff --git a/test/shadowing/alias/let.wgsl.expected.msl b/test/shadowing/alias/let.wgsl.expected.msl
new file mode 100644
index 0000000..292a0ba
--- /dev/null
+++ b/test/shadowing/alias/let.wgsl.expected.msl
@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+void f() {
+  {
+    int const a_1 = int();
+    int const b = a_1;
+  }
+  int const a_2 = int();
+  int const b = a_2;
+}
+
diff --git a/test/shadowing/alias/let.wgsl.expected.spvasm b/test/shadowing/alias/let.wgsl.expected.spvasm
new file mode 100644
index 0000000..56bd05b
--- /dev/null
+++ b/test/shadowing/alias/let.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 9
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %8 = OpConstantNull %int
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %1
+          %6 = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test/shadowing/alias/let.wgsl.expected.wgsl b/test/shadowing/alias/let.wgsl.expected.wgsl
new file mode 100644
index 0000000..558efe6
--- /dev/null
+++ b/test/shadowing/alias/let.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+type a = i32;
+
+fn f() {
+  {
+    let a : a = a();
+    let b = a;
+  }
+  let a : a = a();
+  let b = a;
+}
diff --git a/test/shadowing/alias/param.wgsl b/test/shadowing/alias/param.wgsl
new file mode 100644
index 0000000..42d95bf
--- /dev/null
+++ b/test/shadowing/alias/param.wgsl
@@ -0,0 +1,5 @@
+type a = i32;
+
+fn f(a : a) {
+  let b = a;
+}
diff --git a/test/shadowing/alias/param.wgsl.expected.hlsl b/test/shadowing/alias/param.wgsl.expected.hlsl
new file mode 100644
index 0000000..b48a6b1
--- /dev/null
+++ b/test/shadowing/alias/param.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+void f(int a_1) {
+  const int b = a_1;
+}
diff --git a/test/shadowing/alias/param.wgsl.expected.msl b/test/shadowing/alias/param.wgsl.expected.msl
new file mode 100644
index 0000000..770f18e
--- /dev/null
+++ b/test/shadowing/alias/param.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+void f(int a_1) {
+  int const b = a_1;
+}
+
diff --git a/test/shadowing/alias/param.wgsl.expected.spvasm b/test/shadowing/alias/param.wgsl.expected.spvasm
new file mode 100644
index 0000000..18181e6
--- /dev/null
+++ b/test/shadowing/alias/param.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 10
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %a_1 "a_1"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %5 = OpTypeFunction %void %int
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %5
+        %a_1 = OpFunctionParameter %int
+          %9 = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test/shadowing/alias/param.wgsl.expected.wgsl b/test/shadowing/alias/param.wgsl.expected.wgsl
new file mode 100644
index 0000000..42d95bf
--- /dev/null
+++ b/test/shadowing/alias/param.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+type a = i32;
+
+fn f(a : a) {
+  let b = a;
+}
diff --git a/test/shadowing/alias/var.wgsl b/test/shadowing/alias/var.wgsl
new file mode 100644
index 0000000..10e0c25
--- /dev/null
+++ b/test/shadowing/alias/var.wgsl
@@ -0,0 +1,10 @@
+type a = i32;
+
+fn f() {
+  {
+    var a : a = a();
+    var b = a;
+  }
+  var a : a = a();
+  var b = a;
+}
diff --git a/test/shadowing/alias/var.wgsl.expected.hlsl b/test/shadowing/alias/var.wgsl.expected.hlsl
new file mode 100644
index 0000000..7b1a2a0
--- /dev/null
+++ b/test/shadowing/alias/var.wgsl.expected.hlsl
@@ -0,0 +1,13 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+void f() {
+  {
+    int a_1 = 0;
+    int b = a_1;
+  }
+  int a_2 = 0;
+  int b = a_2;
+}
diff --git a/test/shadowing/alias/var.wgsl.expected.msl b/test/shadowing/alias/var.wgsl.expected.msl
new file mode 100644
index 0000000..9958607
--- /dev/null
+++ b/test/shadowing/alias/var.wgsl.expected.msl
@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+void f() {
+  {
+    int a_1 = int();
+    int b = a_1;
+  }
+  int a_2 = int();
+  int b = a_2;
+}
+
diff --git a/test/shadowing/alias/var.wgsl.expected.spvasm b/test/shadowing/alias/var.wgsl.expected.spvasm
new file mode 100644
index 0000000..0fec746
--- /dev/null
+++ b/test/shadowing/alias/var.wgsl.expected.spvasm
@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %a_1 "a_1"
+               OpName %b "b"
+               OpName %a_2 "a_2"
+               OpName %b_0 "b"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %8 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %1
+          %6 = OpLabel
+        %a_1 = OpVariable %_ptr_Function_int Function %8
+          %b = OpVariable %_ptr_Function_int Function %8
+        %a_2 = OpVariable %_ptr_Function_int Function %8
+        %b_0 = OpVariable %_ptr_Function_int Function %8
+               OpStore %a_1 %8
+         %11 = OpLoad %int %a_1
+               OpStore %b %11
+               OpStore %a_2 %8
+         %14 = OpLoad %int %a_2
+               OpStore %b_0 %14
+               OpReturn
+               OpFunctionEnd
diff --git a/test/shadowing/alias/var.wgsl.expected.wgsl b/test/shadowing/alias/var.wgsl.expected.wgsl
new file mode 100644
index 0000000..10e0c25
--- /dev/null
+++ b/test/shadowing/alias/var.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+type a = i32;
+
+fn f() {
+  {
+    var a : a = a();
+    var b = a;
+  }
+  var a : a = a();
+  var b = a;
+}
diff --git a/test/shadowing/function/let.wgsl b/test/shadowing/function/let.wgsl
new file mode 100644
index 0000000..fb5f04a
--- /dev/null
+++ b/test/shadowing/function/let.wgsl
@@ -0,0 +1,8 @@
+fn a() {
+  {
+    var a = 1;
+    var b = a;
+  }
+  let a = 1;
+  let b = a;
+}
diff --git a/test/shadowing/function/let.wgsl.expected.hlsl b/test/shadowing/function/let.wgsl.expected.hlsl
new file mode 100644
index 0000000..33c805b
--- /dev/null
+++ b/test/shadowing/function/let.wgsl.expected.hlsl
@@ -0,0 +1,12 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+void a() {
+  {
+    int a_1 = 1;
+    int b = a_1;
+  }
+  const int b = 1;
+}
diff --git a/test/shadowing/function/let.wgsl.expected.msl b/test/shadowing/function/let.wgsl.expected.msl
new file mode 100644
index 0000000..16f3280
--- /dev/null
+++ b/test/shadowing/function/let.wgsl.expected.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void a() {
+  {
+    int a_1 = 1;
+    int b = a_1;
+  }
+  int const a_2 = 1;
+  int const b = a_2;
+}
+
diff --git a/test/shadowing/function/let.wgsl.expected.spvasm b/test/shadowing/function/let.wgsl.expected.spvasm
new file mode 100644
index 0000000..592904d
--- /dev/null
+++ b/test/shadowing/function/let.wgsl.expected.spvasm
@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %a "a"
+               OpName %a_1 "a_1"
+               OpName %b "b"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %11 = OpConstantNull %int
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %a = OpFunction %void None %1
+          %6 = OpLabel
+        %a_1 = OpVariable %_ptr_Function_int Function %11
+          %b = OpVariable %_ptr_Function_int Function %11
+               OpStore %a_1 %int_1
+         %12 = OpLoad %int %a_1
+               OpStore %b %12
+               OpReturn
+               OpFunctionEnd
diff --git a/test/shadowing/function/let.wgsl.expected.wgsl b/test/shadowing/function/let.wgsl.expected.wgsl
new file mode 100644
index 0000000..fb5f04a
--- /dev/null
+++ b/test/shadowing/function/let.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+fn a() {
+  {
+    var a = 1;
+    var b = a;
+  }
+  let a = 1;
+  let b = a;
+}
diff --git a/test/shadowing/function/param.wgsl b/test/shadowing/function/param.wgsl
new file mode 100644
index 0000000..9415ae1
--- /dev/null
+++ b/test/shadowing/function/param.wgsl
@@ -0,0 +1,3 @@
+fn a(a : i32) {
+  let b = a;
+}
diff --git a/test/shadowing/function/param.wgsl.expected.hlsl b/test/shadowing/function/param.wgsl.expected.hlsl
new file mode 100644
index 0000000..f06b78f
--- /dev/null
+++ b/test/shadowing/function/param.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+void a(int a_1) {
+  const int b = a_1;
+}
diff --git a/test/shadowing/function/param.wgsl.expected.msl b/test/shadowing/function/param.wgsl.expected.msl
new file mode 100644
index 0000000..1c6280d
--- /dev/null
+++ b/test/shadowing/function/param.wgsl.expected.msl
@@ -0,0 +1,7 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void a(int a_1) {
+  int const b = a_1;
+}
+
diff --git a/test/shadowing/function/param.wgsl.expected.spvasm b/test/shadowing/function/param.wgsl.expected.spvasm
new file mode 100644
index 0000000..de4e783
--- /dev/null
+++ b/test/shadowing/function/param.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 10
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %a "a"
+               OpName %a_1 "a_1"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %5 = OpTypeFunction %void %int
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %a = OpFunction %void None %5
+        %a_1 = OpFunctionParameter %int
+          %9 = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test/shadowing/function/param.wgsl.expected.wgsl b/test/shadowing/function/param.wgsl.expected.wgsl
new file mode 100644
index 0000000..9415ae1
--- /dev/null
+++ b/test/shadowing/function/param.wgsl.expected.wgsl
@@ -0,0 +1,3 @@
+fn a(a : i32) {
+  let b = a;
+}
diff --git a/test/shadowing/function/var.wgsl b/test/shadowing/function/var.wgsl
new file mode 100644
index 0000000..7b3d8ea
--- /dev/null
+++ b/test/shadowing/function/var.wgsl
@@ -0,0 +1,10 @@
+struct a { a : i32; };
+
+fn f() {
+  {
+    var a : a = a();
+    var b = a;
+  }
+  var a : a = a();
+  var b = a;
+}
diff --git a/test/shadowing/function/var.wgsl.expected.hlsl b/test/shadowing/function/var.wgsl.expected.hlsl
new file mode 100644
index 0000000..96ed541
--- /dev/null
+++ b/test/shadowing/function/var.wgsl.expected.hlsl
@@ -0,0 +1,17 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+struct a {
+  int a;
+};
+
+void f() {
+  {
+    a a_1 = (a)0;
+    a b = a_1;
+  }
+  a a_2 = (a)0;
+  a b = a_2;
+}
diff --git a/test/shadowing/function/var.wgsl.expected.msl b/test/shadowing/function/var.wgsl.expected.msl
new file mode 100644
index 0000000..ad6b654
--- /dev/null
+++ b/test/shadowing/function/var.wgsl.expected.msl
@@ -0,0 +1,16 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct a {
+  int a;
+};
+
+void f() {
+  {
+    a a_1 = {};
+    a b = a_1;
+  }
+  a a_2 = {};
+  a b = a_2;
+}
+
diff --git a/test/shadowing/function/var.wgsl.expected.spvasm b/test/shadowing/function/var.wgsl.expected.spvasm
new file mode 100644
index 0000000..2cf9d6b
--- /dev/null
+++ b/test/shadowing/function/var.wgsl.expected.spvasm
@@ -0,0 +1,42 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %a "a"
+               OpMemberName %a 0 "a"
+               OpName %a_1 "a_1"
+               OpName %b "b"
+               OpName %a_2 "a_2"
+               OpName %b_0 "b"
+               OpMemberDecorate %a 0 Offset 0
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %a = OpTypeStruct %int
+          %9 = OpConstantNull %a
+%_ptr_Function_a = OpTypePointer Function %a
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %1
+          %6 = OpLabel
+        %a_1 = OpVariable %_ptr_Function_a Function %9
+          %b = OpVariable %_ptr_Function_a Function %9
+        %a_2 = OpVariable %_ptr_Function_a Function %9
+        %b_0 = OpVariable %_ptr_Function_a Function %9
+               OpStore %a_1 %9
+         %12 = OpLoad %a %a_1
+               OpStore %b %12
+               OpStore %a_2 %9
+         %15 = OpLoad %a %a_2
+               OpStore %b_0 %15
+               OpReturn
+               OpFunctionEnd
diff --git a/test/shadowing/function/var.wgsl.expected.wgsl b/test/shadowing/function/var.wgsl.expected.wgsl
new file mode 100644
index 0000000..feea3e6
--- /dev/null
+++ b/test/shadowing/function/var.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+struct a {
+  a : i32;
+};
+
+fn f() {
+  {
+    var a : a = a();
+    var b = a;
+  }
+  var a : a = a();
+  var b = a;
+}
diff --git a/test/shadowing/param/function.wgsl b/test/shadowing/param/function.wgsl
new file mode 100644
index 0000000..581a6d9
--- /dev/null
+++ b/test/shadowing/param/function.wgsl
@@ -0,0 +1,6 @@
+fn a(a : i32) {
+  {
+    var a = a;
+    var b = a;
+  }
+}
diff --git a/test/shadowing/param/function.wgsl.expected.hlsl b/test/shadowing/param/function.wgsl.expected.hlsl
new file mode 100644
index 0000000..831cce9
--- /dev/null
+++ b/test/shadowing/param/function.wgsl.expected.hlsl
@@ -0,0 +1,11 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+void a(int a_1) {
+  {
+    int a_2 = a_1;
+    int b = a_2;
+  }
+}
diff --git a/test/shadowing/param/function.wgsl.expected.msl b/test/shadowing/param/function.wgsl.expected.msl
new file mode 100644
index 0000000..8be9be4
--- /dev/null
+++ b/test/shadowing/param/function.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void a(int a_1) {
+  {
+    int a_2 = a_1;
+    int b = a_2;
+  }
+}
+
diff --git a/test/shadowing/param/function.wgsl.expected.spvasm b/test/shadowing/param/function.wgsl.expected.spvasm
new file mode 100644
index 0000000..e6291f3
--- /dev/null
+++ b/test/shadowing/param/function.wgsl.expected.spvasm
@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %a "a"
+               OpName %a_1 "a_1"
+               OpName %a_2 "a_2"
+               OpName %b "b"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %5 = OpTypeFunction %void %int
+%_ptr_Function_int = OpTypePointer Function %int
+         %12 = OpConstantNull %int
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %a = OpFunction %void None %5
+        %a_1 = OpFunctionParameter %int
+          %9 = OpLabel
+        %a_2 = OpVariable %_ptr_Function_int Function %12
+          %b = OpVariable %_ptr_Function_int Function %12
+               OpStore %a_2 %a_1
+         %13 = OpLoad %int %a_2
+               OpStore %b %13
+               OpReturn
+               OpFunctionEnd
diff --git a/test/shadowing/param/function.wgsl.expected.wgsl b/test/shadowing/param/function.wgsl.expected.wgsl
new file mode 100644
index 0000000..581a6d9
--- /dev/null
+++ b/test/shadowing/param/function.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+fn a(a : i32) {
+  {
+    var a = a;
+    var b = a;
+  }
+}
diff --git a/test/shadowing/param/let.wgsl b/test/shadowing/param/let.wgsl
new file mode 100644
index 0000000..7a78295
--- /dev/null
+++ b/test/shadowing/param/let.wgsl
@@ -0,0 +1,6 @@
+fn f(a : i32) {
+  {
+    let a = a;
+    let b = a;
+  }
+}
diff --git a/test/shadowing/param/let.wgsl.expected.hlsl b/test/shadowing/param/let.wgsl.expected.hlsl
new file mode 100644
index 0000000..ad522f3
--- /dev/null
+++ b/test/shadowing/param/let.wgsl.expected.hlsl
@@ -0,0 +1,10 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+void f(int a) {
+  {
+    const int b = a;
+  }
+}
diff --git a/test/shadowing/param/let.wgsl.expected.msl b/test/shadowing/param/let.wgsl.expected.msl
new file mode 100644
index 0000000..9f402b1
--- /dev/null
+++ b/test/shadowing/param/let.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void f(int a) {
+  {
+    int const a_1 = a;
+    int const b = a_1;
+  }
+}
+
diff --git a/test/shadowing/param/let.wgsl.expected.spvasm b/test/shadowing/param/let.wgsl.expected.spvasm
new file mode 100644
index 0000000..1cf9296
--- /dev/null
+++ b/test/shadowing/param/let.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 10
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %a "a"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %5 = OpTypeFunction %void %int
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %5
+          %a = OpFunctionParameter %int
+          %9 = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test/shadowing/param/let.wgsl.expected.wgsl b/test/shadowing/param/let.wgsl.expected.wgsl
new file mode 100644
index 0000000..7a78295
--- /dev/null
+++ b/test/shadowing/param/let.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+fn f(a : i32) {
+  {
+    let a = a;
+    let b = a;
+  }
+}
diff --git a/test/shadowing/param/var.wgsl b/test/shadowing/param/var.wgsl
new file mode 100644
index 0000000..e439214
--- /dev/null
+++ b/test/shadowing/param/var.wgsl
@@ -0,0 +1,6 @@
+fn f(a : i32) {
+  {
+    var a = a;
+    var b = a;
+  }
+}
diff --git a/test/shadowing/param/var.wgsl.expected.hlsl b/test/shadowing/param/var.wgsl.expected.hlsl
new file mode 100644
index 0000000..33ab9f1
--- /dev/null
+++ b/test/shadowing/param/var.wgsl.expected.hlsl
@@ -0,0 +1,11 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+void f(int a) {
+  {
+    int a_1 = a;
+    int b = a_1;
+  }
+}
diff --git a/test/shadowing/param/var.wgsl.expected.msl b/test/shadowing/param/var.wgsl.expected.msl
new file mode 100644
index 0000000..a26ff19
--- /dev/null
+++ b/test/shadowing/param/var.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void f(int a) {
+  {
+    int a_1 = a;
+    int b = a_1;
+  }
+}
+
diff --git a/test/shadowing/param/var.wgsl.expected.spvasm b/test/shadowing/param/var.wgsl.expected.spvasm
new file mode 100644
index 0000000..36f0e15
--- /dev/null
+++ b/test/shadowing/param/var.wgsl.expected.spvasm
@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %a "a"
+               OpName %a_1 "a_1"
+               OpName %b "b"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %5 = OpTypeFunction %void %int
+%_ptr_Function_int = OpTypePointer Function %int
+         %12 = OpConstantNull %int
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %5
+          %a = OpFunctionParameter %int
+          %9 = OpLabel
+        %a_1 = OpVariable %_ptr_Function_int Function %12
+          %b = OpVariable %_ptr_Function_int Function %12
+               OpStore %a_1 %a
+         %13 = OpLoad %int %a_1
+               OpStore %b %13
+               OpReturn
+               OpFunctionEnd
diff --git a/test/shadowing/param/var.wgsl.expected.wgsl b/test/shadowing/param/var.wgsl.expected.wgsl
new file mode 100644
index 0000000..e439214
--- /dev/null
+++ b/test/shadowing/param/var.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+fn f(a : i32) {
+  {
+    var a = a;
+    var b = a;
+  }
+}
diff --git a/test/shadowing/struct/let.wgsl b/test/shadowing/struct/let.wgsl
new file mode 100644
index 0000000..5379971
--- /dev/null
+++ b/test/shadowing/struct/let.wgsl
@@ -0,0 +1,10 @@
+struct a { a : i32; };
+
+fn f() {
+  {
+    let a : a = a();
+    let b = a;
+  }
+  let a : a = a();
+  let b = a;
+}
diff --git a/test/shadowing/struct/let.wgsl.expected.hlsl b/test/shadowing/struct/let.wgsl.expected.hlsl
new file mode 100644
index 0000000..c4cb81d
--- /dev/null
+++ b/test/shadowing/struct/let.wgsl.expected.hlsl
@@ -0,0 +1,17 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+struct a {
+  int a;
+};
+
+void f() {
+  {
+    const a a_1 = (a)0;
+    const a b = a_1;
+  }
+  const a a_2 = (a)0;
+  const a b = a_2;
+}
diff --git a/test/shadowing/struct/let.wgsl.expected.msl b/test/shadowing/struct/let.wgsl.expected.msl
new file mode 100644
index 0000000..4780a5e
--- /dev/null
+++ b/test/shadowing/struct/let.wgsl.expected.msl
@@ -0,0 +1,16 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct a {
+  int a;
+};
+
+void f() {
+  {
+    a const a_1 = {};
+    a const b = a_1;
+  }
+  a const a_2 = {};
+  a const b = a_2;
+}
+
diff --git a/test/shadowing/struct/let.wgsl.expected.spvasm b/test/shadowing/struct/let.wgsl.expected.spvasm
new file mode 100644
index 0000000..5119099
--- /dev/null
+++ b/test/shadowing/struct/let.wgsl.expected.spvasm
@@ -0,0 +1,27 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 10
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %a "a"
+               OpMemberName %a 0 "a"
+               OpMemberDecorate %a 0 Offset 0
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %a = OpTypeStruct %int
+          %9 = OpConstantNull %a
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %1
+          %6 = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test/shadowing/struct/let.wgsl.expected.wgsl b/test/shadowing/struct/let.wgsl.expected.wgsl
new file mode 100644
index 0000000..4425a83
--- /dev/null
+++ b/test/shadowing/struct/let.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+struct a {
+  a : i32;
+};
+
+fn f() {
+  {
+    let a : a = a();
+    let b = a;
+  }
+  let a : a = a();
+  let b = a;
+}
diff --git a/test/shadowing/struct/param.wgsl b/test/shadowing/struct/param.wgsl
new file mode 100644
index 0000000..6c79281
--- /dev/null
+++ b/test/shadowing/struct/param.wgsl
@@ -0,0 +1,5 @@
+struct a { a : i32; };
+
+fn f(a : a) {
+  let b = a;
+}
diff --git a/test/shadowing/struct/param.wgsl.expected.hlsl b/test/shadowing/struct/param.wgsl.expected.hlsl
new file mode 100644
index 0000000..03c3814
--- /dev/null
+++ b/test/shadowing/struct/param.wgsl.expected.hlsl
@@ -0,0 +1,12 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+struct a {
+  int a;
+};
+
+void f(a a_1) {
+  const a b = a_1;
+}
diff --git a/test/shadowing/struct/param.wgsl.expected.msl b/test/shadowing/struct/param.wgsl.expected.msl
new file mode 100644
index 0000000..e524263
--- /dev/null
+++ b/test/shadowing/struct/param.wgsl.expected.msl
@@ -0,0 +1,11 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct a {
+  int a;
+};
+
+void f(a a_1) {
+  a const b = a_1;
+}
+
diff --git a/test/shadowing/struct/param.wgsl.expected.spvasm b/test/shadowing/struct/param.wgsl.expected.spvasm
new file mode 100644
index 0000000..ceab72e
--- /dev/null
+++ b/test/shadowing/struct/param.wgsl.expected.spvasm
@@ -0,0 +1,29 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 11
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %a "a"
+               OpMemberName %a 0 "a"
+               OpName %f "f"
+               OpName %a_1 "a_1"
+               OpMemberDecorate %a 0 Offset 0
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %a = OpTypeStruct %int
+          %5 = OpTypeFunction %void %a
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %5
+        %a_1 = OpFunctionParameter %a
+         %10 = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test/shadowing/struct/param.wgsl.expected.wgsl b/test/shadowing/struct/param.wgsl.expected.wgsl
new file mode 100644
index 0000000..2cb3206
--- /dev/null
+++ b/test/shadowing/struct/param.wgsl.expected.wgsl
@@ -0,0 +1,7 @@
+struct a {
+  a : i32;
+};
+
+fn f(a : a) {
+  let b = a;
+}
diff --git a/test/shadowing/struct/var.wgsl b/test/shadowing/struct/var.wgsl
new file mode 100644
index 0000000..7b3d8ea
--- /dev/null
+++ b/test/shadowing/struct/var.wgsl
@@ -0,0 +1,10 @@
+struct a { a : i32; };
+
+fn f() {
+  {
+    var a : a = a();
+    var b = a;
+  }
+  var a : a = a();
+  var b = a;
+}
diff --git a/test/shadowing/struct/var.wgsl.expected.hlsl b/test/shadowing/struct/var.wgsl.expected.hlsl
new file mode 100644
index 0000000..96ed541
--- /dev/null
+++ b/test/shadowing/struct/var.wgsl.expected.hlsl
@@ -0,0 +1,17 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+struct a {
+  int a;
+};
+
+void f() {
+  {
+    a a_1 = (a)0;
+    a b = a_1;
+  }
+  a a_2 = (a)0;
+  a b = a_2;
+}
diff --git a/test/shadowing/struct/var.wgsl.expected.msl b/test/shadowing/struct/var.wgsl.expected.msl
new file mode 100644
index 0000000..ad6b654
--- /dev/null
+++ b/test/shadowing/struct/var.wgsl.expected.msl
@@ -0,0 +1,16 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct a {
+  int a;
+};
+
+void f() {
+  {
+    a a_1 = {};
+    a b = a_1;
+  }
+  a a_2 = {};
+  a b = a_2;
+}
+
diff --git a/test/shadowing/struct/var.wgsl.expected.spvasm b/test/shadowing/struct/var.wgsl.expected.spvasm
new file mode 100644
index 0000000..2cf9d6b
--- /dev/null
+++ b/test/shadowing/struct/var.wgsl.expected.spvasm
@@ -0,0 +1,42 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %a "a"
+               OpMemberName %a 0 "a"
+               OpName %a_1 "a_1"
+               OpName %b "b"
+               OpName %a_2 "a_2"
+               OpName %b_0 "b"
+               OpMemberDecorate %a 0 Offset 0
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %a = OpTypeStruct %int
+          %9 = OpConstantNull %a
+%_ptr_Function_a = OpTypePointer Function %a
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %1
+          %6 = OpLabel
+        %a_1 = OpVariable %_ptr_Function_a Function %9
+          %b = OpVariable %_ptr_Function_a Function %9
+        %a_2 = OpVariable %_ptr_Function_a Function %9
+        %b_0 = OpVariable %_ptr_Function_a Function %9
+               OpStore %a_1 %9
+         %12 = OpLoad %a %a_1
+               OpStore %b %12
+               OpStore %a_2 %9
+         %15 = OpLoad %a %a_2
+               OpStore %b_0 %15
+               OpReturn
+               OpFunctionEnd
diff --git a/test/shadowing/struct/var.wgsl.expected.wgsl b/test/shadowing/struct/var.wgsl.expected.wgsl
new file mode 100644
index 0000000..feea3e6
--- /dev/null
+++ b/test/shadowing/struct/var.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+struct a {
+  a : i32;
+};
+
+fn f() {
+  {
+    var a : a = a();
+    var b = a;
+  }
+  var a : a = a();
+  var b = a;
+}