resolver: Support shadowing

Add transform::Unshadow to renamed shadowed symbols. Required by a
number of other transforms.

Replace Resolver symbol resolution with dep-graph.

The dependency graph now performs full symbol resolution before the
regular resolver pass.
Make use of this instead of duplicating the effort.

Simplfies code, and actually performs variable shadowing consistently.

Fixed: tint:819
Bug: tint:1266
Change-Id: I595d1812aebe1d79d2d32e724ff90de36e74cf4b
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/70523
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

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)
 

src/BUILD.gn

@@ -474,6 +474,8 @@ libtint_source_set("libtint_core_all_src") {
     "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",

src/CMakeLists.txt

@@ -338,6 +338,8 @@ set(TINT_LIB_SRCS
   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 @@ if(${TINT_BUILD_TESTS})
       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

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 @@ Program Parse(const std::vector<uint32_t>& input) {
   // 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;

src/resolver/dependency_graph.cc

@@ -177,6 +177,9 @@ class DependencyScanner {
     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 @@ class DependencyScanner {
       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);

src/resolver/dependency_graph.h

@@ -55,6 +55,12 @@ struct DependencyGraph {
   /// 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

src/resolver/dependency_graph_test.cc

@@ -88,6 +88,17 @@ static constexpr SymbolDeclKind kValueDeclKinds[] = {
     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,
@@ -1171,6 +1182,53 @@ INSTANTIATE_TEST_SUITE_P(Functions,
 
 }  // 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
 ////////////////////////////////////////////////////////////////////////////////

src/resolver/function_validation_test.cc

@@ -26,25 +26,46 @@ namespace {
 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());
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_EQ(r()->error(), "");
 }
 
-TEST_F(ResolverFunctionValidationTest,
+TEST_F(ResolverFunctionValidationTest, ParameterMayShadowGlobal) {
-       ParameterNamesMustBeUniqueShadowsGlobal_pass) {
   // 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());
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_EQ(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 @@ TEST_F(ResolverFunctionValidationTest,
            Decl(var),
        });
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest, FunctionUsingSameVariableName_Pass) {
@@ -74,7 +95,7 @@ TEST_F(ResolverFunctionValidationTest, FunctionUsingSameVariableName_Pass) {
        },
        ast::DecorationList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest,
@@ -95,7 +116,7 @@ TEST_F(ResolverFunctionValidationTest,
        },
        ast::DecorationList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest, UnreachableCode_return) {
@@ -189,7 +210,7 @@ TEST_F(ResolverFunctionValidationTest,
   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 @@ TEST_F(ResolverFunctionValidationTest,
            Return(),
        });
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest,
@@ -269,7 +290,7 @@ TEST_F(ResolverFunctionValidationTest,
        },
        ast::DecorationList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest,
@@ -298,7 +319,7 @@ TEST_F(ResolverFunctionValidationTest,
        },
        ast::DecorationList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest,
@@ -362,7 +383,7 @@ TEST_F(ResolverFunctionValidationTest, NoPipelineEntryPoints) {
        },
        ast::DecorationList{});
 
-  EXPECT_TRUE(r()->Resolve()) << r()->error();
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverFunctionValidationTest, FunctionVarInitWithParam) {
@@ -376,7 +397,7 @@ TEST_F(ResolverFunctionValidationTest, FunctionVarInitWithParam) {
   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 @@ TEST_F(ResolverFunctionValidationTest, FunctionConstInitWithParam) {
   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 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_ConstU32) {
        {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 @@ TEST_F(ResolverFunctionValidationTest, WorkgroupSize_GoodType_U32) {
        {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 @@ TEST_F(ResolverFunctionValidationTest, ParameterSotreType_AtomicFree) {
   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 @@ TEST_F(ResolverFunctionValidationTest, ParametersAtLimit) {
   }
   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 @@ TEST_P(ResolverFunctionParameterValidationTest, StorageClass) {
   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;

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 @@ bool Resolver::ResolveInternal() {
 
   AllocateOverridableConstantIds();
 
+  SetShadows();
+
   if (!ValidatePipelineStages()) {
     return false;
   }
@@ -258,14 +259,8 @@ sem::Type* Resolver::Type(const ast::Type* ty) {
     if (ty->As<ast::ExternalTexture>()) {
       return builder_->create<sem::ExternalTexture>();
     }
-    if (auto* t = ty->As<ast::TypeName>()) {
+    if (auto* type = ResolvedSymbol<sem::Type>(ty)) {
-      auto it = named_type_info_.find(t->name);
+      return type;
-      if (it == named_type_info_.end()) {
-        AddError("unknown type '" + builder_->Symbols().NameFor(t->name) + "'",
-                 t->source);
-        return nullptr;
-      }
-      return it->second.sem;
     }
     TINT_UNREACHABLE(Resolver, diagnostics_)
         << "Unhandled ast::Type: " << ty->TypeInfo().name;
@@ -423,7 +418,7 @@ sem::Variable* Resolver::Variable(const ast::Variable* var,
     }
     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 @@ void Resolver::AllocateOverridableConstantIds() {
   }
 }
 
-bool Resolver::GlobalVariable(const ast::Variable* var) {
+void Resolver::SetShadows() {
-  if (!ValidateNoDuplicateDefinition(var->symbol, var->source,
+  for (auto it : dependencies_.shadows) {
-                                     /* check_global_scope_only */ true)) {
+    auto* var = Sem(it.first);
-    return false;
+    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 @@ bool Resolver::GlobalVariable(const ast::Variable* var) {
 }
 
 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 @@ sem::Function* Resolver::Function(const ast::Function* decl) {
       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 @@ sem::Function* Resolver::Function(const ast::Function* decl) {
     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 @@ sem::Call* Resolver::Call(const ast::CallExpression* expr) {
   auto* ident = expr->target.name;
   Mark(ident);
 
-  auto it = named_type_info_.find(ident->symbol);
+  auto* resolved = ResolvedSymbol(ident);
-  if (it != named_type_info_.end()) {
+  if (auto* ty = As<sem::Type>(resolved)) {
-    // We have a type.
+    return type_ctor_or_conv(ty);
-    return type_ctor_or_conv(it->second.sem);
+  }
+
+  if (auto* fn = As<sem::Function>(resolved)) {
+    return FunctionCall(expr, fn, std::move(args));
   }
 
-  // Not a type, treat as a intrinsic / function call.
   auto name = builder_->Symbols().NameFor(ident->symbol);
   auto intrinsic_type = sem::ParseIntrinsicType(name);
-  auto* call = (intrinsic_type != sem::IntrinsicType::kNone)
+  if (intrinsic_type != sem::IntrinsicType::kNone) {
-                   ? IntrinsicCall(expr, intrinsic_type, std::move(args),
+    return IntrinsicCall(expr, intrinsic_type, std::move(args),
-                                   std::move(arg_tys))
+                         std::move(arg_tys));
-                   : FunctionCall(expr, std::move(args));
+  }
 
-  current_function_->AddDirectCall(call);
+  TINT_ICE(Resolver, diagnostics_)
-  return call;
+      << 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 @@ sem::Call* Resolver::IntrinsicCall(
     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 @@ sem::Call* Resolver::FunctionCall(
     }
   }
 
+  target->AddCallSite(call);
+
   if (!ValidateFunctionCall(call)) {
     return nullptr;
   }
@@ -1476,7 +1466,8 @@ sem::Expression* Resolver::Literal(const ast::LiteralExpression* literal) {
 
 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 @@ sem::Expression* Resolver::Identifier(const ast::IdentifierExpression* expr) {
     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 (IsIntrinsic(symbol)) {
-  if (sem::ParseIntrinsicType(name) != sem::IntrinsicType::kNone) {
     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 @@ sem::Expression* Resolver::UnaryOp(const ast::UnaryOpExpression* unary) {
 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 @@ bool Resolver::VariableDeclStatement(const ast::VariableDeclStatement* stmt) {
     }
   }
 
-  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 @@ sem::Type* Resolver::TypeDecl(const ast::TypeDecl* named_type) {
     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 @@ sem::Array* Resolver::Array(const ast::Array* arr) {
 
     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 @@ bool Resolver::Scope(sem::CompoundStatement* stmt, F&& callback) {
   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 @@ bool Resolver::IsHostShareable(const sem::Type* type) const {
   return false;
 }
 
+bool Resolver::IsIntrinsic(Symbol symbol) const {
+  std::string name = builder_->Symbols().NameFor(symbol);
+  return sem::ParseIntrinsicType(name) != sem::IntrinsicType::kNone;
+}
+
 ////////////////////////////////////////////////////////////////////////////////
 // Resolver::TypeConversionSig
 ////////////////////////////////////////////////////////////////////////////////

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 @@ class Resolver {
   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 @@ class Resolver {
   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 @@ class Resolver {
                                        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 @@ class Resolver {
   /// 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 @@ class Resolver {
   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 @@ class Resolver {
   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,

src/resolver/resolver_validation.cc

@@ -942,11 +942,6 @@ bool Resolver::ValidateInterpolateDecoration(
 
 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 @@ bool Resolver::ValidateFunctionCall(const sem::Call* call) {
     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 @@ bool Resolver::ValidateFunctionCall(const sem::Call* call) {
         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 @@ bool Resolver::ValidateScalarConstructorOrCast(const ast::CallExpression* ctor,
   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,24 +2316,23 @@ bool Resolver::ValidateAssignment(const ast::AssignmentStatement* a) {
   // 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 = ResolvedSymbol<sem::Variable>(a->lhs)) {
-    if (auto* var = variable_stack_.Get(ident->symbol)) {
+    auto* decl = var->Declaration();
     if (var->Is<sem::Parameter>()) {
       AddError("cannot assign to function parameter", a->lhs->source);
-        AddNote("'" + builder_->Symbols().NameFor(ident->symbol) +
+      AddNote("'" + builder_->Symbols().NameFor(decl->symbol) +
                   "' is declared here:",
-                var->Declaration()->source);
+              decl->source);
       return false;
     }
-      if (var->Declaration()->is_const) {
+    if (decl->is_const) {
       AddError("cannot assign to const", a->lhs->source);
-        AddNote("'" + builder_->Symbols().NameFor(ident->symbol) +
+      AddNote("'" + builder_->Symbols().NameFor(decl->symbol) +
                   "' is declared here:",
-                var->Declaration()->source);
+              decl->source);
       return false;
     }
   }
-  }
 
   auto* lhs_ref = lhs_ty->As<sem::Reference>();
   if (!lhs_ref) {
@@ -2388,36 +2365,6 @@ bool Resolver::ValidateAssignment(const ast::AssignmentStatement* a) {
   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;

src/resolver/validation_test.cc

@@ -150,25 +150,28 @@ TEST_F(ResolverValidationTest, Expr_ErrUnknownExprType) {
 
 TEST_F(ResolverValidationTest, Expr_DontCall_Function) {
   Func("func", {}, ty.void_(), {}, {});
-  auto* ident = create<ast::IdentifierExpression>(
+  WrapInFunction(Expr(Source{{{3, 3}, {3, 8}}}, "func"));
-      Source{{Source::Location{3, 3}, Source::Location{3, 8}}},
-      Symbols().Register("func"));
-  WrapInFunction(ident);
 
   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>(
+  WrapInFunction(Expr(Source{{{3, 3}, {3, 8}}}, "round"));
-      Source{{Source::Location{3, 3}, Source::Location{3, 8}}},
-      Symbols().Register("round"));
-  WrapInFunction(ident);
 
   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 @@ TEST_F(ResolverValidationTest, UsingUndefinedVariableGlobalVariable_Pass) {
 
   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 @@ TEST_F(ResolverValidationTest, UsingUndefinedVariableInnerScope_Fail) {
   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, StorageClass_TextureExplicitStorageClass) {
 TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadChar) {
   Global("my_vec", ty.vec3<f32>(), ast::StorageClass::kPrivate);
 
-  auto* ident = create<ast::IdentifierExpression>(
+  auto* ident = Expr(Source{{{3, 3}, {3, 7}}}, "xyqz");
-      Source{{Source::Location{3, 3}, Source::Location{3, 7}}},
-      Symbols().Register("xyqz"));
 
   auto* mem = MemberAccessor("my_vec", ident);
   WrapInFunction(mem);
@@ -364,9 +365,7 @@ TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadChar) {
 TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_MixedChars) {
   Global("my_vec", ty.vec4<f32>(), ast::StorageClass::kPrivate);
 
-  auto* ident = create<ast::IdentifierExpression>(
+  auto* ident = Expr(Source{{{3, 3}, {3, 7}}}, "rgyw");
-      Source{{Source::Location{3, 3}, Source::Location{3, 7}}},
-      Symbols().Register("rgyw"));
 
   auto* mem = MemberAccessor("my_vec", ident);
   WrapInFunction(mem);
@@ -380,9 +379,7 @@ TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_MixedChars) {
 TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadLength) {
   Global("my_vec", ty.vec3<f32>(), ast::StorageClass::kPrivate);
 
-  auto* ident = create<ast::IdentifierExpression>(
+  auto* ident = Expr(Source{{{3, 3}, {3, 8}}}, "zzzzz");
-      Source{{Source::Location{3, 3}, Source::Location{3, 8}}},
-      Symbols().Register("zzzzz"));
   auto* mem = MemberAccessor("my_vec", ident);
   WrapInFunction(mem);
 
@@ -393,8 +390,7 @@ TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadLength) {
 TEST_F(ResolverValidationTest, Expr_MemberAccessor_VectorSwizzle_BadIndex) {
   Global("my_vec", ty.vec2<f32>(), ast::StorageClass::kPrivate);
 
-  auto* ident = create<ast::IdentifierExpression>(Source{{3, 3}},
+  auto* ident = Expr(Source{{3, 3}}, "z");
-                                                  Symbols().Register("z"));
   auto* mem = MemberAccessor("my_vec", ident);
   WrapInFunction(mem);
 
@@ -406,9 +402,7 @@ TEST_F(ResolverValidationTest, Expr_MemberAccessor_BadParent) {
   // var param: vec4<f32>
   // let ret: f32 = *(&param).x;
   auto* param = Var("param", ty.vec4<f32>());
-  auto* x = create<ast::IdentifierExpression>(
+  auto* x = Expr(Source{{{3, 3}, {3, 8}}}, "x");
-      Source{{Source::Location{3, 3}, Source::Location{3, 8}}},
-      Symbols().Register("x"));
 
   auto* addressOf_expr = AddressOf(Source{{12, 34}}, param);
   auto* accessor_expr = MemberAccessor(addressOf_expr, x);
@@ -430,9 +424,7 @@ TEST_F(ResolverValidationTest, EXpr_MemberAccessor_FuncGoodParent) {
   auto* p =
       Param("p", ty.pointer(ty.vec4<f32>(), ast::StorageClass::kFunction));
   auto* star_p = Deref(p);
-  auto* z = create<ast::IdentifierExpression>(
+  auto* z = Expr(Source{{{3, 3}, {3, 8}}}, "z");
-      Source{{Source::Location{3, 3}, Source::Location{3, 8}}},
-      Symbols().Register("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 @@ TEST_F(ResolverValidationTest, EXpr_MemberAccessor_FuncBadParent) {
   // }
   auto* p =
       Param("p", ty.pointer(ty.vec4<f32>(), ast::StorageClass::kFunction));
-  auto* z = create<ast::IdentifierExpression>(
+  auto* z = Expr(Source{{{3, 3}, {3, 8}}}, "z");
-      Source{{Source::Location{3, 3}, Source::Location{3, 8}}},
-      Symbols().Register("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 @@ TEST_F(ResolverValidationTest,
   //     }
   // }
 
-  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 @@ TEST_F(ResolverValidationTest,
   //     }
   // }
 
-  auto cont_loc = Source{Source::Location{12, 34}};
+  auto cont_loc = Source{{12, 34}};
-  auto decl_loc = Source{Source::Location{56, 78}};
+  auto decl_loc = Source{{56, 78}};
-  auto ref_loc = Source{Source::Location{90, 12}};
+  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 @@ TEST_F(
   //     }
   // }
 
-  auto cont_loc = Source{Source::Location{12, 34}};
+  auto cont_loc = Source{{12, 34}};
-  auto decl_loc = Source{Source::Location{56, 78}};
+  auto decl_loc = Source{{56, 78}};
-  auto ref_loc = Source{Source::Location{90, 12}};
+  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 @@ TEST_F(ResolverValidationTest,
   //     }
   // }
 
-  auto cont_loc = Source{Source::Location{12, 34}};
+  auto cont_loc = Source{{12, 34}};
-  auto decl_loc = Source{Source::Location{56, 78}};
+  auto decl_loc = Source{{56, 78}};
-  auto ref_loc = Source{Source::Location{90, 12}};
+  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 @@ TEST_F(ResolverValidationTest,
   //     }
   // }
 
-  auto cont_loc = Source{Source::Location{12, 34}};
+  auto cont_loc = Source{{12, 34}};
-  auto decl_loc = Source{Source::Location{56, 78}};
+  auto decl_loc = Source{{56, 78}};
-  auto ref_loc = Source{Source::Location{90, 12}};
+  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 @@ TEST_F(ResolverTest, Stmt_Loop_ContinueInLoopBodyAfterDecl_UsageInContinuing) {
   //     }
   // }
 
-  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));

src/resolver/var_let_test.cc

@@ -58,7 +58,7 @@ TEST_F(ResolverVarLetTest, TypeOfVar) {
            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 @@ TEST_F(ResolverVarLetTest, TypeOfLet) {
            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 @@ TEST_F(ResolverVarLetTest, DefaultVarStorageClass) {
 
   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 @@ TEST_F(ResolverVarLetTest, ExplicitVarStorageClass) {
                              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 @@ TEST_F(ResolverVarLetTest, LetInheritsAccessFromOriginatingVariable) {
 
   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 @@ TEST_F(ResolverVarLetTest, LetInheritsAccessFromOriginatingVariable) {
   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

src/resolver/var_let_validation_test.cc

@@ -175,10 +175,7 @@ TEST_F(ResolverVarLetValidationTest, GlobalVarRedeclaredAsLocal) {
   WrapInFunction(Var(Source{{12, 34}}, "v", ty.f32(), ast::StorageClass::kNone,
                      Expr(2.0f)));
 
-  EXPECT_FALSE(r()->Resolve()) << r()->error();
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_EQ(
-      r()->error(),
-      "12:34 error: redefinition of 'v'\nnote: previous definition is here");
 }
 
 TEST_F(ResolverVarLetValidationTest, VarRedeclaredInInnerBlock) {
@@ -194,10 +191,7 @@ TEST_F(ResolverVarLetValidationTest, VarRedeclaredInInnerBlock) {
 
   WrapInFunction(outer_body);
 
-  EXPECT_FALSE(r()->Resolve());
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_EQ(
-      r()->error(),
-      "12:34 error: redefinition of 'v'\nnote: previous definition is here");
 }
 
 TEST_F(ResolverVarLetValidationTest, VarRedeclaredInIfBlock) {
@@ -219,10 +213,7 @@ TEST_F(ResolverVarLetValidationTest, VarRedeclaredInIfBlock) {
 
   WrapInFunction(outer_body);
 
-  EXPECT_FALSE(r()->Resolve());
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
-  EXPECT_EQ(
-      r()->error(),
-      "12:34 error: redefinition of 'v'\nnote: previous definition is here");
 }
 
 TEST_F(ResolverVarLetValidationTest, InferredPtrStorageAccessMismatch) {

src/sem/type_mappings.h

@@ -25,6 +25,7 @@ class CallExpression;
 class Expression;
 class Function;
 class MemberAccessorExpression;
+class Node;
 class Statement;
 class Struct;
 class StructMember;
@@ -56,11 +57,12 @@ struct TypeMappings {
   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
 };

src/sem/variable.cc

@@ -45,12 +45,10 @@ LocalVariable::LocalVariable(const ast::Variable* declaration,
                              const sem::Type* type,
                              ast::StorageClass storage_class,
                              ast::Access access,
+                             const sem::Statement* statement,
                              Constant constant_value)
-    : Base(declaration,
+    : Base(declaration, type, storage_class, access, std::move(constant_value)),
-           type,
+      statement_(statement) {}
-           storage_class,
-           access,
-           std::move(constant_value)) {}
 
 LocalVariable::~LocalVariable() = default;
 

src/sem/variable.h

@@ -95,15 +95,31 @@ class LocalVariable : public Castable<LocalVariable, Variable> {
   /// @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 @@ class Parameter : public Castable<Parameter, Variable> {
   /// @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

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 @@ TEST_F(ArrayLengthFromUniformTest, Error_MissingTransformData) {
       "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 @@ fn main() {
   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 @@ fn main() {
   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 @@ fn main() {
   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 @@ fn main() {
   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 @@ fn main() {
   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 @@ fn main() {
   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}),

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 @@ CalculateArrayLength::~CalculateArrayLength() = default;
 
 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

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 @@ namespace {
 
 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 @@ fn main() {
 }
 )";
 
-  auto got = Run<CalculateArrayLength>(src);
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -105,7 +119,7 @@ fn main() {
 }
 )";
 
-  auto got = Run<CalculateArrayLength>(src);
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -149,7 +163,7 @@ fn main() {
 }
 )";
 
-  auto got = Run<CalculateArrayLength>(src);
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -206,7 +220,7 @@ fn main() {
 }
 )";
 
-  auto got = Run<CalculateArrayLength>(src);
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -274,7 +288,63 @@ fn main() {
 }
 )";
 
-  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));
 }

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 @@ struct CanonicalizeEntryPointIO::State {
 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(

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 @@ namespace {
 
 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 @@ TEST_F(CanonicalizeEntryPointIOTest, Error_MissingTransformData) {
       "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 @@ fn comp_main() {
   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 @@ fn frag_main() {
   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 @@ fn frag_main([[builtin(position)]] coord : vec4<f32>, tint_symbol : tint_symbol_
   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 @@ fn frag_main(tint_symbol : tint_symbol_1) {
   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 @@ fn frag_main(tint_symbol : tint_symbol_1) {
   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 @@ fn frag_main() {
   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 @@ fn frag_main([[builtin(position)]] coord : vec4<f32>, tint_symbol : tint_symbol_
   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 @@ fn frag_main(tint_symbol : tint_symbol_1) {
   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 @@ fn frag_main() {
   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 @@ fn frag_main() -> tint_symbol {
   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 @@ fn frag_main() -> tint_symbol {
   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 @@ fn frag_main() {
   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 @@ fn frag_main() -> tint_symbol {
   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 @@ fn frag_main() -> tint_symbol {
   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 @@ fn frag_main2() {
   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 @@ fn frag_main2(tint_symbol_2 : tint_symbol_3) {
   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 @@ fn frag_main2(tint_symbol_2 : tint_symbol_3) {
   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 @@ fn frag_main1(tint_symbol : tint_symbol_1) {
   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 @@ fn frag_main(tint_symbol : tint_symbol_1) -> tint_symbol_2 {
   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 @@ fn frag_main(tint_symbol_1 : tint_symbol_2) {
   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 @@ fn frag_main() {
   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 @@ fn main2() -> tint_symbol_1 {
   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 @@ fn frag_main(tint_symbol : tint_symbol_1) -> tint_symbol_2 {
   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 @@ fn frag_main(tint_symbol_1 : tint_symbol_2) {
   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 @@ fn tint_symbol_1(tint_symbol : tint_symbol_2) {
   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 @@ fn frag_main() -> tint_symbol {
   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 @@ fn frag_main() -> tint_symbol {
   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 @@ fn frag_main() -> tint_symbol {
   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 @@ fn frag_main() -> tint_symbol {
   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 @@ fn frag_main() -> tint_symbol {
   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 @@ fn frag_main() -> tint_symbol {
   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 @@ fn comp_main1() {
   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 @@ fn frag_main() -> tint_symbol {
   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 @@ fn vert_main() {
   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 @@ fn vert_main() -> tint_symbol {
   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 @@ fn vert_main() {
   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 @@ fn vert_main() -> tint_symbol {
   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 @@ fn vert_main() {
   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 @@ fn vert_main(tint_symbol : tint_symbol_1) -> tint_symbol_2 {
   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 @@ fn vert_main(tint_symbol : tint_symbol_1) -> tint_symbol_2 {
   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 @@ fn main() {
   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));
 }

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 @@ TEST_F(DecomposeStridedMatrixTest, Empty) {
   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 @@ fn f() {
 }
 )";
 
-  auto got =
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -171,8 +172,8 @@ fn f() {
 }
 )";
 
-  auto got =
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -234,8 +235,8 @@ fn f() {
 }
 )";
 
-  auto got =
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -300,8 +301,8 @@ fn f() {
 }
 )";
 
-  auto got =
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -362,8 +363,8 @@ fn f() {
 }
 )";
 
-  auto got =
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -429,8 +430,8 @@ fn f() {
 }
 )";
 
-  auto got =
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -491,8 +492,8 @@ fn f() {
 }
 )";
 
-  auto got =
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -580,8 +581,8 @@ fn f() {
 }
 )";
 
-  auto got =
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -637,8 +638,8 @@ fn f() {
 }
 )";
 
-  auto got =
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }
@@ -695,8 +696,8 @@ fn f() {
 }
 )";
 
-  auto got =
+  auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedMatrix>(
-      Run<SimplifyPointers, DecomposeStridedMatrix>(Program(std::move(b)));
+      Program(std::move(b)));
 
   EXPECT_EQ(expect, str(got));
 }

src/transform/fold_trivial_single_use_lets_test.cc

@@ -35,13 +35,13 @@ TEST_F(FoldTrivialSingleUseLetsTest, Single) {
   auto* src = R"(
 fn f() {
   let x = 1;
-  ignore(x);
+  _ = x;
 }
 )";
 
   auto* expect = R"(
 fn f() {
-  ignore(1);
+  _ = 1;
 }
 )";
 
@@ -56,13 +56,13 @@ fn f() {
   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 @@ fn f() {
   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 function_with_posssible_side_effect() -> i32 {
 fn f() {
   let x = 1;
   let y = function_with_posssible_side_effect();
-  ignore((x + y));
+  _ = (x + y);
 }
 )";
 
@@ -117,7 +117,7 @@ TEST_F(FoldTrivialSingleUseLetsTest, NoFold_UseInSubBlock) {
 fn f() {
   let x = 1;
   {
-    ignore(x);
+    _ = x;
   }
 }
 )";
@@ -133,7 +133,26 @@ TEST_F(FoldTrivialSingleUseLetsTest, NoFold_MultipleUses) {
   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);
+  }
 }
 )";
 

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 @@ Output Glsl::Run(const Program* in, const DataMap& inputs) {
 
   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>();

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 @@ TEST_F(NumWorkgroupsFromUniformTest, Error_MissingTransformData) {
   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 @@ fn main() {
   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 @@ fn main() {
   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 @@ fn main(tint_symbol : tint_symbol_1) {
   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 @@ fn main3() {
   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 @@ fn main(tint_symbol : tint_symbol_1) {
   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));
 }
 

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;
 SimplifyPointers::~SimplifyPointers() = default;
 
 void SimplifyPointers::Run(CloneContext& ctx, const DataMap&, DataMap&) {
+  if (!Requires<Unshadow>(ctx)) {
+    return;
+  }
   State(ctx).Run();
 }
 

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.
 ///

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 @@ namespace {
 
 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 @@ fn f() {
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -79,7 +88,7 @@ fn f() {
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -103,7 +112,7 @@ fn f() {
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -126,7 +135,7 @@ fn f() {
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -180,7 +189,7 @@ fn matrix() {
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -201,7 +210,7 @@ fn f() {
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -229,7 +238,7 @@ fn f() {
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -264,7 +273,7 @@ fn main() {
 }
 )";
 
-  auto got = Run<SimplifyPointers>(src);
+  auto got = Run<Unshadow, SimplifyPointers>(src);
 
   EXPECT_EQ(expect, str(got));
 }
@@ -330,34 +339,42 @@ fn f() {
 }
 )";
 
-  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, ShadowPointer) {
-TEST_F(SimplifyPointersTest, DISABLED_ModificationAfterInline) {
   auto* src = R"(
-fn x(p : ptr<function, i32>) -> i32 {
+var<private> a : array<i32, 2>;
-  return *p;
-}
 
-fn f() {
+[[stage(compute), workgroup_size(1)]]
-  var i : i32 = 1;
+fn main() {
-  let p : ptr<function, i32> = &i;
+  let x = &a;
-  if (true) {
+  var a : i32 = (*x)[0];
-    var i : i32 = 2;
+  {
-    x(p);
+    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

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

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_

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

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 @@ SanitizedResult Sanitize(
   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>();

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 @@ SanitizedResult Sanitize(
       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.

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 @@ SanitizedResult Sanitize(const Program* in,
   transform::Manager manager;
   transform::DataMap data;
 
+  manager.Add<transform::Unshadow>();
   if (!disable_workgroup_init) {
     manager.Add<transform::ZeroInitWorkgroupMemory>();
   }

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));
+}

test/BUILD.gn

@@ -327,6 +327,7 @@ tint_unittests_source_set("tint_unittests_transform_src") {
     "../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",

test/bug/tint/825.wgsl.expected.hlsl

@@ -7,5 +7,5 @@ void f() {
   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];
 }

test/bug/tint/825.wgsl.expected.msl

@@ -5,6 +5,6 @@ void f() {
   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];
 }
 

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;
+}

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;
+}

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;
+}
+

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

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;
+}

test/shadowing/alias/param.wgsl

@@ -0,0 +1,5 @@
+type a = i32;
+
+fn f(a : a) {
+  let b = a;
+}

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;
+}

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;
+}
+

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

test/shadowing/alias/param.wgsl.expected.wgsl

@@ -0,0 +1,5 @@
+type a = i32;
+
+fn f(a : a) {
+  let b = a;
+}

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;
+}

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;
+}

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;
+}
+

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

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;
+}

test/shadowing/function/let.wgsl

@@ -0,0 +1,8 @@
+fn a() {
+  {
+    var a = 1;
+    var b = a;
+  }
+  let a = 1;
+  let b = a;
+}

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;
+}

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;
+}
+

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

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;
+}

test/shadowing/function/param.wgsl

@@ -0,0 +1,3 @@
+fn a(a : i32) {
+  let b = a;
+}

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;
+}

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;
+}
+

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

test/shadowing/function/param.wgsl.expected.wgsl

@@ -0,0 +1,3 @@
+fn a(a : i32) {
+  let b = a;
+}

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;
+}

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;
+}

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;
+}
+

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

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;
+}

test/shadowing/param/function.wgsl

@@ -0,0 +1,6 @@
+fn a(a : i32) {
+  {
+    var a = a;
+    var b = a;
+  }
+}

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;
+  }
+}

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;
+  }
+}
+

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

test/shadowing/param/function.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+fn a(a : i32) {
+  {
+    var a = a;
+    var b = a;
+  }
+}

test/shadowing/param/let.wgsl

@@ -0,0 +1,6 @@
+fn f(a : i32) {
+  {
+    let a = a;
+    let b = a;
+  }
+}

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;
+  }
+}

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;
+  }
+}
+

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

test/shadowing/param/let.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+fn f(a : i32) {
+  {
+    let a = a;
+    let b = a;
+  }
+}

test/shadowing/param/var.wgsl

@@ -0,0 +1,6 @@
+fn f(a : i32) {
+  {
+    var a = a;
+    var b = a;
+  }
+}

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;
+  }
+}

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;
+  }
+}
+

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

test/shadowing/param/var.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+fn f(a : i32) {
+  {
+    var a = a;
+    var b = a;
+  }
+}

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;
+}

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;
+}

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;
+}
+

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

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;
+}

test/shadowing/struct/param.wgsl

@@ -0,0 +1,5 @@
+struct a { a : i32; };
+
+fn f(a : a) {
+  let b = a;
+}

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;
+}

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;
+}
+

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

test/shadowing/struct/param.wgsl.expected.wgsl

@@ -0,0 +1,7 @@
+struct a {
+  a : i32;
+};
+
+fn f(a : a) {
+  let b = a;
+}

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;
+}

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;
+}

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;
+}
+

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

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;
+}