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